ó
ŒúçLc           @   s–  d  Z  d d l m Z d Z d Z d d l m Z d d l Z d d l m Z m	 Z	 m
 Z
 d e g  d	 d	 d	 d
 „ Z d e g  d „ Z d e g  d „ Z d d0 d „  ƒ  YZ e d d „ ƒ Z e d d „ ƒ Z e d d „ ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z d „  Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d  „  ƒ Z  e d! „  ƒ Z! e d" „  ƒ Z" e d# „  ƒ Z# e d$ „  ƒ Z$ e d% „  ƒ Z% e d& „  ƒ Z& e d' „  ƒ Z' e d d( „ ƒ Z) e d) „  ƒ Z* e d* „  ƒ Z+ e d+ „  ƒ Z, e d, „  ƒ Z- e d- „  ƒ Z. e d. „  ƒ Z/ e d/ „  ƒ Z0 d S(1   s   Elementwise functionality.iÿÿÿÿ(   t   divisions$   Copyright (C) 2009 Andreas Kloeckners   
Permission is hereby granted, free of charge, to any person
obtaining a copy of this software and associated documentation
files (the "Software"), to deal in the Software without
restriction, including without limitation the rights to use,
copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the
Software is furnished to do so, subject to the following
conditions:

The above copyright notice and this permission notice shall be
included in all copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
OTHER DEALINGS IN THE SOFTWARE.
(   t   context_dependent_memoizeN(   t   dtype_to_ctypet	   VectorArgt	   ScalarArgt   kernelt    c   	      C   sg   d d l  m } | d i d j d „  |  Dƒ ƒ d 6| d 6| d 6| d	 6| d
 6| d 6d | d | ƒS(   Niÿÿÿÿ(   t   SourceModulesÞ  
        #include <pycuda-complex.hpp>

        %(preamble)s

        __global__ void %(name)s(%(arguments)s)
        {

          unsigned tid = threadIdx.x;
          unsigned total_threads = gridDim.x*blockDim.x;
          unsigned cta_start = blockDim.x*blockIdx.x;
          unsigned i;

          %(loop_prep)s;

          for (i = cta_start + tid; i < n; i += total_threads)
          {
            %(operation)s;
          }

          %(after_loop)s;
        }
        s   , c         s   s   |  ] } | j  ƒ  Vq d  S(   N(   t
   declarator(   t   .0t   arg(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>C   s    t	   argumentst	   operationt   namet   preamblet	   loop_prept
   after_loopt   optionst   keep(   t   pycuda.compilerR   t   join(	   R   R   R   R   R   R   R   R   R   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_elwise_module(   s    c         K   sÊ   t  |  t ƒ rJ d d l m } g  |  j d ƒ D] } | | ƒ ^ q/ }  n  |  j t t j d ƒ ƒ t	 |  | | | | |  } d d l m
 }	 | j | ƒ }
 |
 j d j d „  |  Dƒ ƒ d	 ƒ |
 |  f S(
   Niÿÿÿÿ(   t   parse_c_argt   ,t   n(   t   get_arg_typeR   c         s   s   |  ] } | j  Vq d  S(   N(   t   struct_char(   R	   R
   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>Y   s    i   (   i   i   i   (   t
   isinstancet   strt   pycuda.toolsR   t   splitt   appendR   t   numpyt   uintpR   R   t   get_functiont   prepareR   (   R   R   R   R   R   t   kwargsR   R
   t   modR   t   func(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_elwise_kernel_and_typesL   s    +#c         K   s%   t  |  | | | | |  \ } }  | S(   sn   Return a L{pycuda.driver.Function} that performs the same scalar operation
    on one or several vectors.
    (   R'   (   R   R   R   R   R   R$   R&   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_elwise_kernel]   s    t   ElementwiseKernelc           B   s#   e  Z d  e g  d „ Z d „  Z RS(   R   c   	      K   sk   t  | | | | | |  \ |  _ |  _ g  t |  j ƒ D]! \ } } t | t ƒ r7 | ^ q7 sg t d ‚ d  S(   NsX   ElementwiseKernel can only be used with functions that have at least one vector argument(   R'   R&   R   t	   enumerateR   R   t   AssertionError(	   t   selfR   R   R   R   R   R$   t   iR
   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   __init__k   s
    $c         G   s²   g  } g  } x\ t  | |  j ƒ D]H \ } } t | t ƒ rZ | j | ƒ | j | j ƒ q | j | ƒ q W| d } | j | j ƒ |  j j | j	 Œ  |  j j
 | j | Œ d  S(   Ni    (   t   zipR   R   R   R   t   gpudatat   mem_sizeR&   t   set_block_shapet   _blockt   prepared_callt   _grid(   R,   t   argst   vectorst   invocation_argsR
   t	   arg_descrt   repr_vec(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   __call__u   s    
(   t   __name__t
   __module__t   FalseR.   R;   (    (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyR)   j   s   	i   c   
         sf  i t  | ƒ d 6t  |  ƒ d 6t  |  d t ƒd 6‰  t | d ƒ g g  t | ƒ D] } t |  d t | ƒ ƒ ^ qO t t j d ƒ g } d d	 j ‡  f d
 †  t | ƒ Dƒ ƒ } d ˆ  d	 j d „  t | ƒ Dƒ ƒ } t	 | | d d | ƒ} | j
 d ƒ } g  t | ƒ D] } | j d | ƒ ^ q}	 | j d | d t j t j ƒ j d d |	 ƒ| |	 f S(   Nt   idx_tpt   tpt   with_fp_tex_hackt   tex_tpt   idxt   destR   s   #include <pycuda-helpers.hpp>

s   
c         3   s#   |  ] } d  ˆ  d | f Vq d S(   s3   texture <%s, 1, cudaReadModeElementType> tex_src%d;RB   N(    (   R	   R-   (   t   ctx(    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>”   s   s   %(idx_tp)s src_idx = idx[i];
c         s   s   |  ] } d  | | f Vq d S(   s.   dest%d[i] = fp_tex1Dfetch(tex_src%d, src_idx);N(    (   R	   R-   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>™   s   t   takeR   s	   tex_src%dt   Pi   t   texrefs(   i   i   i   (   R   t   TrueR   t   rangeR   R   R    t   intpR   R   R"   t
   get_texrefR#   t   dtypeR!   t   char(
   RM   t	   idx_dtypet	   vec_countR-   R6   R   t   bodyR%   R&   t   tex_src(    (   RE   s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_take_kernelˆ   s$    ?	,1c            sÒ  i t  | ƒ d 6t  |  ƒ d 6t  |  d t ƒd 6‰ t | d ƒ t | d ƒ g g  t | ƒ D] } t |  d | ƒ ^ q[ g  t | ƒ D] } | r… t | d | ƒ ^ q… t t j d	 ƒ g } d
 d j ‡ f d †  t | ƒ Dƒ ƒ } | rù d „  ‰  n	 d „  ‰  d ˆ d j ‡  f d †  t | ƒ Dƒ ƒ } t | | d d | ƒ} | j	 d ƒ }	 g  t | ƒ D] } | j
 d | ƒ ^ qc}
 |	 j d | d t | ƒ | | j t j t j ƒ j d d |
 ƒ|	 |
 f S(   NR?   R@   RA   RB   t   gmem_dest_idxt   gmem_src_idxs   dest%ds   offset%dR   s   #include <pycuda-helpers.hpp>

s   
c         3   s#   |  ] } d  ˆ  d | f Vq d S(   s3   texture <%s, 1, cudaReadModeElementType> tex_src%d;RB   N(    (   R	   R-   (   RE   (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>¹   s   c         S   s   d |  |  |  f S(   Ns>   dest%d[dest_idx] = fp_tex1Dfetch(tex_src%d, src_idx+offset%d);(    (   R-   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_copy_insn½   s    c         S   s   d |  |  f S(   Ns5   dest%d[dest_idx] = fp_tex1Dfetch(tex_src%d, src_idx);(    (   R-   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyRV   Â   s    sN   %(idx_tp)s src_idx = gmem_src_idx[i];
%(idx_tp)s dest_idx = gmem_dest_idx[i];
c         3   s   |  ] } ˆ  | ƒ Vq d  S(   N(    (   R	   R-   (   RV   (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>È   s    t   take_putR   s	   tex_src%dt   PPRG   i   RH   (   i   i   i   (   R   RI   R   RJ   R   R    RK   R   R   R"   RL   R#   t   boolRN   RM   R!   (   RM   RO   t   with_offsetsRP   R-   R6   R   RQ   R%   R&   RR   (    (   RV   RE   s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_take_put_kernel¥   s*    ‹		&,3c         C   s  i t  | ƒ d 6t  |  ƒ d 6} t | d ƒ g g  t | ƒ D] } t |  d | ƒ ^ q< g  t | ƒ D] } t |  d | ƒ ^ qf t t j d ƒ g } d | d j d	 „  t | ƒ Dƒ ƒ } t | | d
 ƒ j d
 ƒ } | j	 d d | d t j
 t j ƒ j d ƒ | S(   NR?   R@   RT   s   dest%ds   src%dR   s(   %(idx_tp)s dest_idx = gmem_dest_idx[i];
s   
c         s   s   |  ] } d  | | f Vq d S(   s   dest%d[dest_idx] = src%d[i];N(    (   R	   R-   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>ë   s   t   putRG   i   i   (   i   i   i   (   R   R   RJ   R   R    RK   R   R   R"   R#   RM   R!   RN   (   RM   RO   RP   RE   R-   R6   RQ   R&   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_put_kernelØ   s    y/c         C   s.   t  d i t |  ƒ d 6t | ƒ d 6d d ƒ S(   Ns"   %(tp_dest)s *dest, %(tp_src)s *srct   tp_destt   tp_srcs   dest[i] = src[i]t   copy(   R(   R   (   t
   dtype_destt	   dtype_src(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_copy_kernelõ   s    c      	   C   s  d d l  m } d d l m } m } m } g  } d g } g  } g  }	 g  }
 xå t |  ƒ D]× \ } \ } } } | ræ | j d | | d t ƒ| f ƒ | j | | d | ƒ ƒ |
 j d | ƒ | j d	 | | ƒ | | f ƒ n4 | j | | d
 | ƒ ƒ | j | | d | ƒ ƒ |	 j d | | f ƒ qZ W| j | | d ƒ ƒ | j | t	 j
 d ƒ ƒ | | d d j |	 ƒ d d d j | ƒ d d j | ƒ ƒ} | j d ƒ } g  |
 D] } | j | ƒ ^ q·} | j d j d „  | Dƒ ƒ d d | ƒ| | f S(   Niÿÿÿÿ(   R   (   R   R   R   s   #include <pycuda-helpers.hpp>

s1   texture <%s, 1, cudaReadModeElementType> tex_a%d;RA   s   x%ds   tex_a%ds"   %s a%d = fp_tex1Dfetch(tex_a%d, 0)s   a%ds
   a%d*x%d[i]t   zR   s   z[i] = s    + t   linear_combinationR   s   
R   s   ;
R   c         s   s   |  ] } | j  Vq d  S(   N(   R   (   R	   R
   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>*  s    i   RH   (   i   i   i   (   R   R   t   pycuda.elementwiseR   R   R   R*   R   RI   R    R!   R   R"   RL   R#   (   t   summand_descriptorst   dtype_zR   R   R   R   R6   R   R   t   summandst	   tex_namesR-   t   is_gpu_scalart   scalar_dtypet   vector_dtypeR%   R&   t   tnRR   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_linear_combination_kernel  sB    	"c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns=   %(tp_x)s a, %(tp_x)s *x, %(tp_y)s b, %(tp_y)s *y, %(tp_z)s *zt   tp_xt   tp_yt   tp_zs   z[i] = a*x[i] + b*y[i]t   axpbyz(   R(   R   (   t   dtype_xt   dtype_yRh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_axpbyz_kernel2  s    c         C   s!   t  d i t |  ƒ d 6d d ƒ S(   Ns'   %(tp)s a, %(tp)s *x,%(tp)s b, %(tp)s *zR@   s   z[i] = a * x[i] + bt   axpb(   R(   R   (   RM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_axpbz_kernel=  s
    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] * y[i]t   multiply(   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_multiply_kernelE  s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] / y[i]t   divide(   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_divide_kernelP  s    c         C   s!   t  d i t |  ƒ d 6d d ƒ S(   Ns   %(tp)s *x, %(tp)s y, %(tp)s *zR@   s   z[i] = y / x[i]t   divide_r(   R(   R   (   RM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_rdivide_elwise_kernel[  s
    c         C   sC   t  d i t | ƒ d 6t | ƒ d 6t | ƒ d 6d |  |  d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = %s(x[i], y[i])t   _kernel(   R(   R   (   R&   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_binary_func_kerneld  s    c         C   sq   t  j | | g k r" |  d }  n  d d l m } | d „  | | | g Dƒ ƒ r^ d |  }  n  t |  | | | ƒ S(   Nt   fiÿÿÿÿ(   t   anyc         s   s   |  ] } | j  d  k Vq d S(   R   N(   t   kind(   R	   t   dt(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pys	   <genexpr>t  s    (   R    t   float64t   pytoolsR‚   R€   (   R&   Rt   Ru   Rh   R‚   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_binary_minmax_kernelo  s    c         C   s!   t  d i t |  ƒ d 6d d ƒ S(   Ns   %(tp)s a, %(tp)s *zR@   s   z[i] = at   fill(   R(   R   (   RM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_fill_kernely  s
    c         C   s!   t  d i t |  ƒ d 6d d ƒ S(   Ns   %(tp)s *y, %(tp)s *zR@   s   z[i] = y[n-1-i]t   reverse(   R(   R   (   RM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_reverse_kernel‚  s
    c         C   s.   t  d i t |  ƒ d 6t | ƒ d 6d d ƒ S(   Ns   %(tp)s *y, %(real_tp)s *zR@   t   real_tps   z[i] = real(y[i])t   real(   R(   R   (   RM   t
   real_dtype(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_real_kernel‹  s    c         C   s.   t  d i t |  ƒ d 6t | ƒ d 6d d ƒ S(   Ns   %(tp)s *y, %(real_tp)s *zR@   RŒ   s   z[i] = imag(y[i])t   imag(   R(   R   (   RM   RŽ   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_imag_kernel•  s    c         C   s!   t  d i t |  ƒ d 6d d ƒ S(   Ns   %(tp)s *y, %(tp)s *zR@   s   z[i] = pycuda::conj(y[i])t   conj(   R(   R   (   RM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_conj_kernelŸ  s
    c         C   s!   t  d i t |  ƒ d 6d d ƒ S(   Ns$   %(tp)s *z, %(tp)s start, %(tp)s stepR@   s   z[i] = start + i*stept   arange(   R(   R   (   RM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_arange_kernel¨  s
    c         C   sC   |  t  j k r d } n d } t d i t |  ƒ d 6d | d ƒ S(   Nt   powt   powfs"   %(tp)s value, %(tp)s *y, %(tp)s *zR@   s   z[i] = %s(y[i], value)t
   pow_method(   R    R…   R(   R   (   RM   R&   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_pow_kernel²  s    	c         C   sc   t  j |  | g k r d } n d } t d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d | d ƒ S(	   NR–   R—   s%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = %s(x[i], y[i])R˜   (   R    R…   R(   R   (   Rt   Ru   Rh   R&   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_pow_array_kernelÀ  s    	c           C   s   t  d d d ƒ S(   Ns    float *arg, float *mod, float *zs   z[i] = fmod(arg[i], mod[i])t   fmod_kernel(   R(   (    (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_fmod_kernelÐ  s    c           C   s   t  d d d ƒ S(   Ns)   float *x, float *intpart ,float *fracparts%   fracpart[i] = modf(x[i], &intpart[i])t   modf_kernel(   R(   (    (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_modf_kernel×  s    c           C   s   t  d d d ƒ S(   Ns-   float *x, float *significand, float *exponents„   
                int expt = 0;
                significand[i] = frexp(x[i], &expt);
                exponent[i] = expt;
            t   frexp_kernel(   R(   (    (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_frexp_kernelÞ  s    c           C   s   t  d d d ƒ S(   Ns!   float *sig, float *expt, float *zs"   z[i] = ldexp(sig[i], int(expt[i]))t   ldexp_kernel(   R(   (    (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_ldexp_kernelé  s    c         C   sK   | d  k r | } n  t d i t | ƒ d 6t | ƒ d 6d |  d |  ƒ S(   Ns   %(tp_in)s *y, %(tp_out)s *zt   tp_int   tp_outs   z[i] = %s(y[i])s	   %s_kernel(   t   NoneR(   R   (   t	   func_namet   in_dtypet	   out_dtype(    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_unary_func_kernelð  s    	c         C   s@   t  t |  d ƒ t | d ƒ t | d ƒ t | d ƒ g d d ƒ S(   Nt   critt   then_t   else_t   results-   result[i] = crit[i] > 0 ? then_[i] : else_[i]t   if_positive(   R(   R   (   t
   crit_dtypeRM   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_if_positive_kernel   s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] == y[i]t   eq(   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_eq_kernel  s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] != y[i]t   ne(   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_ne_kernel  s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] <= y[i]t   lt(   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_le_kernel$  s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] >= y[i]Rµ   (   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_ge_kernel/  s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] < y[i]Rµ   (   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_lt_kernel:  s    c         C   s;   t  d i t |  ƒ d 6t | ƒ d 6t | ƒ d 6d d ƒ S(   Ns%   %(tp_x)s *x, %(tp_y)s *y, %(tp_z)s *zRp   Rq   Rr   s   z[i] = x[i] > y[i]t   gt(   R(   R   (   Rt   Ru   Rh   (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   get_gt_kernelE  s    (    (1   t   __doc__t
   __future__R    t   __copyright__t   __license__R   R   R    R   R   R   R>   R   R'   R(   R)   RS   R[   R]   Rc   Ro   Rv   Rx   Rz   R|   R~   R€   R‡   R‰   R‹   R   R‘   R“   R•   R™   Rš   Rœ   Rž   R    R¢   R¥   R©   R°   R²   R´   R¶   R·   R¸   Rº   (    (    (    s3   C:\Python27\Lib\site-packages\pycuda\elementwise.pyt   <module>   s^   	#21		
		

	
