ó
.™xTc           @   s»   d  d l  m Z d  d l Z d  d l Z d  d l m Z d  d l m Z d  d l m	 Z	 d f  d „  ƒ  YZ
 d Z d	 e Z d
 f  d „  ƒ  YZ d f  d „  ƒ  YZ d e f d „  ƒ  YZ d S(   iÿÿÿÿ(   t   divisionN(   t   flex(   t   ext(   t   Sorryt   pprocessc           B   sM   e  Z d  Z d „  Z e d „ Z e d „ Z d „  Z d „  Z e d „ Z	 RS(   s�  Class to represent a worker process

  Software requirements for running direct summation using Cuda parallel processing units:

    1. numpy package (http://numpy.scipy.org) version 1.5.1 OK
    2. pycuda package, version 2011.1 OK

       * more information
          - http://wiki.tiker.net/PyCuda/Installation/Linux
          - http://documen.tician.de/pycuda/tutorial.html
       * suggested install scheme for Linux (not necessarily exact)::

          ../base/bin/python configure.py --cuda-root=/usr/common/usg/cuda/3.2 \
          --cudadrv-lib-dir=/usr/common/usg/nvidia-driver-util/3.2/lib64 \
          --boost-inc-dir=$HOME/boostbuild/include \
          --boost-lib-dir=$HOME/boostbuild/lib \
          --boost-python-libname=boost_python \
          --boost-thread-libname=boost_thread
          make install

    3. gcc 4.4.2 or higher is required for Linux build of pycuda 2011.1
    4. boost.python is required for pycuda; cctbx-installed version is probably OK, not tested.
       tests were performed with separately-installed boost::

          cd boost_1_45_0
          ./bootstrap.sh --prefix=$HOME/boostbuild --libdir=$HOME/boostbuild/lib \
          --with-python=$HOME/build/base/bin/python \
          --with-libraries=signals,thread,python
          ./bjam variant=release link=shared install

    5. cuda 3.2 separately installed is required for pycuda 2011.1
       (http://developer.nvidia.com/object/cuda_3_2_downloads.html)

  Suggested hardware as tested:
    * Nvidia Tesla C2050 (Fermi, compute capability 2.0): 225-fold performance improvement over CPU.
    * Nvidia Tesla C1060 (compute capability 1.3): 24-fold performance improvement.
  c   
      C   se  x& |  j  D] } | GH| j GH| j GHq
 W|  j j ƒ  d } |  j j | ƒ } |  j j |  j  ƒ } xû t t |  j  ƒ ƒ D]ä } |  j  | GH|  j  | j	 GH| | | GH|  j j
 |  j  | j	 ƒ j ƒ  } d G| GH|  j j
 |  j  | j	 ƒ j ƒ  } | d } t j | ƒ }	 xH t d t | ƒ d ƒ D]. } |	 | | t j | | d | ƒ 7}	 q&W|	 GHqy Wd  S(   Ng      ð?t   constantg      @i    i   i   (   t
   scattererst   fpt   fdpt   registryt   showt    unique_form_factors_at_d_star_sqt   unique_indicest   xranget   lent   scattering_typet   gaussiant   ct
   parameterst   copyt   matht   exp(
   t   selft	   scatterert	   d_star_sqt   uniquefft   uniqueixt   xt   constant_termR   t   x_sqt   unique_form_factor_at_x_sq(    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   print_diagnostics2   s(    "	"
,c         C   sî   t  |  j d ƒ r< |  j j ƒ  j ƒ  j ƒ  j | j ƒ } n! |  j j ƒ  j ƒ  j | j ƒ } |  j j ƒ  t d k rÍ d t d |  j j ƒ  t } | rº d | j	 | j
 d | f GHn  | j | f ƒ n  | j
 d d |  _ | |  _ d S(   s�    The number of miller indices and atoms each must be an exact
        multiple of the BLOCKSIZE (32), so zero-padding is employedt   as_vec3_doublei    i   i   s5   Miller indices: resetting %s array from size %d to %dN(   t   hasattrt   miller_indicesR    t	   as_doublet   as_numpy_arrayt   astypet   numpy_tt   sizet   FHKL_BLOCKSIZEt   dtypet   shapet   resizet
   n_flat_hklt   flat_mix(   R   t	   algorithmt   verboseR-   t   newsize(    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   prepare_miller_arrays_for_cudaH   s    	*! c         C   sr  |  j  j d ƒ } t j | ƒ |  j _ t |  j j ƒ |  j _ g  |  j _ g  |  j _	 |  j  j
 |  j ƒ } xO t |  j j ƒ D]; } |  j j j | j | ƒ ƒ |  j j	 j d d g ƒ q Wx± t |  j j ƒ D]� } |  j j	 |  j j | d |  j j | |  j j	 |  j j | d <| d |  j j k  rÑ |  j j	 |  j j | d |  j j	 |  j j | d d <qÑ qÑ W| r­t |  j j ƒ GH|  j j GH|  j j GH|  j j	 GHn  t j | ƒ } | j | ƒ j ƒ  j | j j ƒ } |  j j ƒ  } | j ƒ  }	 | j | ƒ j ƒ  j ƒ  j | j ƒ }
 t j g  |  j D] } | j ƒ  ^ q2ƒ j | ƒ j ƒ  j | j ƒ } |  j j ƒ  j | ƒ j ƒ  j | j ƒ } |	 t d k rAt d |	 t } | rúd |
 j |
 j d d | f GHd | j | j d | f GHHn  |
 j d | f ƒ | j | f ƒ | j | f ƒ | j | f ƒ n  |	 |  _ |
 |  _  | |  _! | |  _" | |  _# d  S(   Ng        i    i   s5   Scatterer xyzs: resetting %s array from size %d to %di   s1   weights   : resetting %s array from size %d to %d($   R	   R   R   t   sort_permutationR   t   increasing_orderR   t   number_of_typest   unsorted_countst   sorted_rangesR   R   t   appendt   countt   listt   selectR$   R%   t   numpyt   uint32t   extract_sitesR'   R#   R&   t   doublet   weightt   extract_u_isoR(   R)   R*   R+   t   n_sitest
   flat_sitest   weightst   u_isoR   (   R   R.   R/   t   n_electronsR   t   st   uniqueix_sort_ordert   sorted_uniqueixt   sitesRA   t   sorted_flat_sitest   St   sorted_weightst   sorted_u_isoR0   (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt!   prepare_scattering_sites_for_cuda^   s\    )A'+				c         C   sI  t  j ƒ  } d |  _ d |  _ xÃ |  j j ƒ  D]² } |  j d 7_ | j | j ƒ  ƒ | j ƒ  } |  j d k r¯ t	 | ƒ d |  j k s� t
 ‚ t	 | ƒ d d k s¯ t
 ‚ n  t	 | ƒ d |  _ x | D] } | j | ƒ qÉ Wq. W| j ƒ  j | j ƒ |  _ d d |  j |  _ |  j j d |  j |  j k s9t
 ‚ |  j j ƒ  |  _ |  j j ƒ  |  _ d |  _ t  j ƒ  } xe |  j D]Z } x' | j ƒ  j ƒ  D] } | j | ƒ q•Wx' | j ƒ  j ƒ  D] } | j | ƒ q¿Wq|W| j ƒ  j | j ƒ |  _ |  j j d |  j |  j k st
 ‚ | j j |  j j ƒ  ƒ j | j ƒ |  _ d  S(   Ni    i   i   i   (   R   R>   t   n_gaussianst   n_terms_in_sumR	   t   unique_gaussians_as_listR7   R   R   R   t   AssertionErrorR$   R%   R&   t	   gaussianst   g_strideR*   t   space_groupt   order_zt   order_pt
   sym_stridet   rR#   t   tt   symmetryR;   t   arrayt	   unit_cellt   reciprocal_metrical_matrix(   R   R.   RS   R   t   termst   itemR[   t   symop(    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt!   prepare_gaussians_symmetries_cell¦   s:    		&	  &	c         C   s   d } d } d } y t  |  j ƒ d k s0 t ‚ Wn t k
 rP t d ƒ ‚ n X|  j j ƒ  j d ƒ t  |  j ƒ k  |  _ y( |  j j	 ƒ  j t
 ƒ d k s¢ t ‚ Wn t k
 rÂ t d ƒ ‚ n Xy~ g  |  j D] } | j ^ qÐ j d ƒ t  |  j ƒ k st ‚ g  |  j D] } | j ^ qj d ƒ t  |  j ƒ k s@t ‚ Wn t k
 r`t d ƒ ‚ n Xy+ | j d ƒ j ƒ  } | d k s‹t ‚ Wn! t k
 r¯t d | ƒ ‚ n Xy( | j d k r×| d k s×t ‚ n  Wn! t k
 rût d | ƒ ‚ n Xt  |  j ƒ }	 |	 t d k r0t d	 |	 t }	 n  t  |  j ƒ }
 |
 t d k rdt d	 |
 t }
 n  d |	 } d |
 } | j d ƒ j ƒ  } d | | | k r½t d | | | f ƒ ‚ n  | j d ƒ j | j j ƒ | j d ƒ j | j j ƒ } |
 | k rt d |
 | f ƒ ‚ n  d  S(   Ns-   Fhkl = Sum(atoms) (weight * exp(2*pi*i(H*X)))so   Fhkl = Sum(atoms(X)) (f0(d*) * weight *
        Sum(symops(SR,ST)) exp(2*pi*i(H*SR*X+H*ST) + DW-factor)
      )sy   Fhkl = Sum(atoms(X)) ((f0(d*)+f'+i*f") * weight *
        Sum(symops(SR,ST)) exp(2*pi*i(H*SR*X+H*ST) + DW-factor)
      )i    s'   There must be at least one Miller indexg        sm   As presently implemented parallel processor direct summation doesn't support anisotropic displacement factorssa   As presently implemented parallel processor direct summation doesn't support anomalous scatterersi   i   sB   Implementation assumes CUDA compute capability >= 1.3; found %d.%dt   floati   sT   Float32 kernel tests correctly only with CUDA compute capability >= 2.0; found %d.%di$   i(   g333333ó?sL   Atoms use %d bytes; hkl indices use %d bytes, exceeding total GPU memory %d.si   %d reflection values will exceed the maximum number of
        threads for a single kernel invocation, %d(   i   i   (   i   i    (   R   R"   RR   t	   ExceptionR   R   R@   R8   t   use_debye_wallert   extract_use_u_anisot   TrueR   R   t   Devicet   compute_capabilityt   float_tR(   t   total_memoryt   get_attributet   device_attributet   MAX_BLOCK_DIM_Xt   MAX_GRID_DIM_X(   R   t   managert   cudaR.   t   PyNX_implementationt   Present_implementationt   Future_implementationRK   t   cct   n_atomst   n_hklt   global_memory_atomst   global_memory_hklt   totalmemt   totalthreads(    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   validate_the_inputsÑ   s^    *(=A 

c         C   sK  | j  j ƒ  |  _ | j  j ƒ  |  _ | j j ƒ  |  _ | j j ƒ  |  _ | j j ƒ  j	 ƒ  |  _ | rs |  j
 ƒ  n  t | d ƒ r¹ t j |  j | j j ƒ  |  j |  j |  j ƒ | _ d St | d ƒ rGd d l j } d d l m } |  j | | | ƒ |  j | ƒ |  j | ƒ |  j | ƒ | j j ƒ  d k s<t ‚ | j d ƒ } | j | j j ƒ } | j | j j ƒ } | t |  j j  d |  j! j  d |  j" |  j# t$ |  j% ƒ |  j& |  j' f j( d | j) ƒ ƒ }	 |	 j* d	 ƒ d }
 | j+ |
 |  j, ƒ |	 j* d
 ƒ d } | j+ | |  j ƒ |	 j* d ƒ d } | j+ | |  j! ƒ |	 j- d ƒ } | j. j/ |  j0 f | j1 ƒ } | j. j/ |  j0 f | j1 ƒ } x:t2 |  j j3 ƒ D]&} | j. j/ |  j0 f | j1 ƒ } | j. j/ |  j0 f | j1 ƒ } | | j4 | ƒ | j4 | ƒ | j5 |  j6 ƒ | j5 |  j7 ƒ | j5 |  j8 ƒ | j. j9 |  j j: | ƒ | j. j9 |  j j; | d ƒ | j. j9 |  j j; | d ƒ | j5 |  j< ƒ d t= d d f d |  j0 t= d f ƒ	| | 7} | | 7} q¢Wt> j? | d t@ |  j ƒ !jA | j. jB ƒ ƒ } t> j? | d t@ |  j ƒ !jA | j. jB ƒ ƒ } tC t> jD | | ƒ ƒ | _ d Sd S(   sÑ  Initializes a direct parallel worker

    :param instance: an instance of class from_scatterers_direct(cctbx.xray.structure_factors.manager.managed_calculation_base)
    :type instance: cctbx.xray.structure_factors.from_scatterers_direct
    :param algorithm: an instance of class direct_summation_cuda_platform(direct_summation_simple) with algorithm set to "simple" or "pycuda"
    :type algorithm: cctbx.xray.structure_factors.direct_summation_cuda_platform
    t   simpleNt   pycudaiÿÿÿÿ(   t   SourceModulei   i    t   floating_point_tR^   RS   R[   t	   CUDA_fhklt   blockt   grid(E   t   _xray_structureR   t   scattering_type_registryR	   t   _miller_sett   indicesR"   R]   RU   t	   make_tidyR   R!   R   t   structure_factors_simplet   _resultst   pycuda.drivert   drivert   pycuda.compilerR   R|   R1   RN   Rb   Rh   R8   RR   Rl   Rm   t	   WARP_SIZEt   MULTIPROCESSOR_COUNTt   mod_fhkl_sortedRS   R*   R[   RX   RT   t   intRe   RV   RW   t   replaceRj   t
   get_globalt   memcpy_htodR^   t   get_functionR;   t   zerosR,   R&   R   R4   t   InOutt   InRB   RC   RD   R<   R3   R6   R-   R(   R   R>   R   R%   t   float64t   fcalc_containert   complex_double(   R   t   instanceR.   R/   Rq   R   t   devicet   WARPSIZER�   t   sort_modt   r_m_m_addresst   gaussian_addresst   symmetry_addressR�   t   intermediate_realt   intermediate_imagR   t	   fhkl_realt	   fhkl_imagt   flex_fhkl_realt   flex_fhkl_imag(    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   __init__  sz    
  %
..(
   t   __name__t
   __module__t   __doc__R   t   FalseR1   RN   Rb   R|   R©   (    (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyR      s   %	H	+	Ii    sí  
__device__ __constant__ floating_point_t reciprocal_metrical_matrix[6];
__device__ __constant__ floating_point_t gaussians[%%d];
__device__ __constant__ floating_point_t symmetry[%%d];

__device__ floating_point_t get_d_star_sq(floating_point_t const& H,floating_point_t const& K,floating_point_t const& L){
   return
            (H * H) * reciprocal_metrical_matrix[0]
          + (K * K) * reciprocal_metrical_matrix[1]
          + (L * L) * reciprocal_metrical_matrix[2]
          + (2 * H * K) * reciprocal_metrical_matrix[3]
          + (2 * H * L) * reciprocal_metrical_matrix[4]
          + (2 * K * L) * reciprocal_metrical_matrix[5];
}

# define SYM_STRIDE %%d

__device__ floating_point_t get_H_dot_translation(
                              floating_point_t const& H,floating_point_t const& K,floating_point_t const& L,
                              unsigned int const & isym){
  // case for primitive cells: return 0
  // but in general,
     return H * (symmetry[SYM_STRIDE*isym+9])
          + K * (symmetry[SYM_STRIDE*isym+10])
          + L * (symmetry[SYM_STRIDE*isym+11])
    ;
}

__device__ floating_point_t get_H_dot_X(floating_point_t* const& site, unsigned int const& i,
                              floating_point_t const& H,floating_point_t const& K,floating_point_t const& L,
                              unsigned int const & isym){
  // case for space group P1
  // return H * site[0] + K * site[1] + L * site[2];

  // but in general,
     return site[3*i+0] * (H * symmetry[SYM_STRIDE*isym]+
                       K * symmetry[SYM_STRIDE*isym+3]+
                       L * symmetry[SYM_STRIDE*isym+6])
          + site[3*i+1] * (H * symmetry[SYM_STRIDE*isym+1]+
                       K * symmetry[SYM_STRIDE*isym+4]+
                       L * symmetry[SYM_STRIDE*isym+7])
          + site[3*i+2] * (H * symmetry[SYM_STRIDE*isym+2]+
                       K * symmetry[SYM_STRIDE*isym+5]+
                       L * symmetry[SYM_STRIDE*isym+8])
          + H * symmetry[SYM_STRIDE*isym+9]
          + K * symmetry[SYM_STRIDE*isym+10]
          + L * symmetry[SYM_STRIDE*isym+11]
    ;
}

__global__ void CUDA_fhkl(floating_point_t *fhkl_real,floating_point_t *fhkl_imag,
                        const floating_point_t *xyzsites, const floating_point_t *weight,
                        const floating_point_t *u_iso,
                        const unsigned int scattering_type,
                        const unsigned int index_begin,
                        const unsigned int index_end,
                        const floating_point_t *hkl)
{
   #define BLOCKSIZE %d
   #define G_STRIDE %%d
   #define USE_DEBYE_WALLER %%d
   #define ORDER_Z %%d
   #define ORDER_P %%d
   #define ORDER_T ORDER_Z/ORDER_P

   const unsigned long ix=threadIdx.x+blockDim.x*blockIdx.x;

   /*later figure out how to make this shared*/
   const floating_point_t twopi= 8.* atan(1.);
   #if USE_DEBYE_WALLER
     const floating_point_t eight_pi_sq = 2.* twopi * twopi;
     __shared__ floating_point_t b_iso[BLOCKSIZE];
   #endif
   /* can the threads coalesce with a stride of 3? Make into __shared__ and See Programming Guide Fig. G-2 */
   const floating_point_t h = hkl[3*ix];
   const floating_point_t k = hkl[3*ix+1];
   const floating_point_t l = hkl[3*ix+2];
   floating_point_t fr=0,fi=0;

   __shared__ floating_point_t xyz[3*BLOCKSIZE];
   __shared__ floating_point_t wght[BLOCKSIZE];

   long atom_p = BLOCKSIZE*index_begin/BLOCKSIZE;
   floating_point_t x_sq = get_d_star_sq(h,k,l) / 4.;

   /* get unique form factors at d_star_sq.
      all scatterers are the same for any particular kernel invocation; hence only one form factor
    */
   floating_point_t form_factor = gaussians[G_STRIDE*scattering_type];// constant term
   for (unsigned int term=1; term < G_STRIDE; term+=2){
     form_factor += gaussians[G_STRIDE*scattering_type + term] *
                    exp(-gaussians[G_STRIDE*scattering_type + term + 1] * x_sq);
   }

   for ( ;atom_p<index_end; atom_p+=BLOCKSIZE ) {
      xyz[threadIdx.x]=xyzsites[3*atom_p+threadIdx.x];
      xyz[BLOCKSIZE+threadIdx.x]=xyzsites[3*atom_p+threadIdx.x+BLOCKSIZE];
      xyz[2*BLOCKSIZE+threadIdx.x]=xyzsites[3*atom_p+threadIdx.x+2*BLOCKSIZE];
      wght[threadIdx.x]=weight[atom_p+threadIdx.x];
      #if USE_DEBYE_WALLER
        b_iso[threadIdx.x] = eight_pi_sq * u_iso[atom_p+threadIdx.x];
      #endif
      __syncthreads();

      for(unsigned int i=0;i<BLOCKSIZE;i++) {
        floating_point_t form_factor_mask = form_factor;
        floating_point_t debye_waller_factor = 1.;
        #if USE_DEBYE_WALLER
          debye_waller_factor *= exp( -b_iso[i] * x_sq );
        #endif

        floating_point_t s,c,s_imag=0.,c_real=0.;
        for(unsigned int isym=0; isym < ORDER_P; isym++){
          /*do not use the faster, but less accurate intrinsic function.*/
          sincos(twopi* get_H_dot_X( xyz, i, h, k, l, isym), &s,&c);
          c_real += c;
          s_imag += s;
        }
        if (atom_p+i >= index_begin && atom_p+i < index_end) {
          fr += form_factor_mask * wght[i] * c_real * debye_waller_factor;
          fi += form_factor_mask * wght[i] * s_imag * debye_waller_factor;
        }
      }
   }

   //multiply factor from centring translations
   floating_point_t s,c,s_imag_trans=0.,c_real_trans=0.;
   for(unsigned int isym=0; isym < ORDER_Z; isym+=ORDER_P){
     sincos(twopi* get_H_dot_translation( h, k, l, isym), &s,&c);
     c_real_trans += c;
     s_imag_trans += s;
   }

   fhkl_real[ix] += (c_real_trans * fr - s_imag_trans * fi);
   fhkl_imag[ix] += (s_imag_trans * fr + c_real_trans * fi);
}
Rš   c           B   s   e  Z d  „  Z d „  Z RS(   c         C   s   | |  _  d  S(   N(   t   fcalc(   R   R®   (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyR©     s    c         C   s   |  j  S(   N(   R®   (   R   (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   f_calc  s    (   Rª   R«   R©   R¯   (    (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyRš      s   	t   direct_summation_simplec           B   s    e  Z e Z d  „  Z d „  Z RS(   c         C   s   t  |  _ d  S(   N(   Rg   R}   (   R   (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyR©     s    c         C   s
   | d k S(   Nt   direct(    (   R   t   other(    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   __eq__	  s    (   Rª   R«   Rg   t   use_alt_parallelR©   R³   (    (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyR°     s   	t   direct_summation_cuda_platformc           B   s   e  Z d  d „ Z d „  Z RS(   R>   c         C   s    t  |  _ | |  _ |  j ƒ  d  S(   N(   Rg   R~   Rj   t   validate_platform_resources(   R   Rj   (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyR©     s    		c         C   s¦   y d d  l  } | |  _  Wn t k
 r8 t d ƒ ‚ n X|  j d k sN t ‚ i | j d 6| j d 6|  j |  _ y d d  l } Wn t k
 r¡ t d ƒ ‚ n Xd  S(   NiÿÿÿÿsG   Module numpy must be installed for parallel processor direct summation.R>   Rc   sH   Module pycuda must be installed for parallel processor direct summation.(   s   doubles   float(	   R;   Rd   R   Rj   RR   R™   t   float32R&   t   pycuda.autoinit(   R   R;   R~   (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyR¶     s    $(   Rª   R«   R©   R¶   (    (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyRµ     s   (   t
   __future__R    R   R   t   scitbx.array_familyR   t
   cctbx.xrayR   t   libtbx.utilsR   R   R(   R�   Rš   R°   Rµ   (    (    (    sj   /public2/ftp/pub/ccp4/6.5/ccp4-6.5/lib/py2/cctbx/xray/structure_factors/from_scatterers_direct_parallel.pyt   <module>   s   ÿ k‰
