Module Expr Str FunctionDef arguments args NameParam NameParam NameParam defaults body Expr Str Assign NameStore Call AttributeLoad Str attr NameLoad FunctionDef arguments args defaults body For NameStore NameLoad body If Call NameLoad NameLoad TupleLoad NameLoad NameLoad body For NameStore NameLoad body Expr Yield NameLoad orelse Expr Yield NameLoad decorator_list Assign NameStore Call AttributeLoad Str attr Call NameLoad Return BinOpMod Str Call NameLoad decorator_list FunctionDef arguments args NameParam defaults body Expr Str If UnaryOpNot Call NameLoad NameLoad NameLoad body Raise Call NameLoad Str NameLoad FunctionDef arguments args NameParam defaults body Assign AttributeStore NameLoad attr NameLoad Return NameLoad decorator_list Return NameLoad decorator_list Assign NameStore TupleLoad FunctionDef arguments args NameParam NameParam NameParam NameParam NameParam defaults body Expr Str Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Return BinOpMod Str Call NameLoad decorator_list Call NameLoad TupleLoad Num FunctionDef arguments args NameParam NameParam NameParam NameParam defaults body Return Call NameLoad NameLoad NameLoad NameLoad NameLoad Lambda arguments args NameParam NameParam defaults BinOpMod Str TupleLoad NameLoad NameLoad decorator_list Call NameLoad AttributeLoad NameLoad attr FunctionDef arguments args NameParam NameParam NameParam NameParam defaults body Return Call NameLoad NameLoad NameLoad NameLoad NameLoad Lambda arguments args NameParam NameParam defaults BinOpMod Str TupleLoad NameLoad NameLoad decorator_list Call NameLoad AttributeLoad NameLoad attr FunctionDef arguments args NameParam NameParam NameParam NameParam defaults body Return Call NameLoad NameLoad NameLoad NameLoad NameLoad Lambda arguments args NameParam NameParam defaults BinOpMod Str TupleLoad NameLoad NameLoad decorator_list Call NameLoad AttributeLoad NameLoad attr FunctionDef arguments args NameParam NameParam NameParam NameParam defaults body Return Call NameLoad NameLoad NameLoad NameLoad NameLoad Lambda arguments args NameParam NameParam defaults BinOpMod Str TupleLoad NameLoad NameLoad decorator_list Call NameLoad AttributeLoad NameLoad attr FunctionDef arguments args NameParam NameParam NameParam NameParam NameParam defaults body Expr Str Return ListLoad Call NameLoad NameLoad NameLoad NameLoad NameLoad Str BinOpAdd BinOpAdd Str NameLoad Str Str BinOpAdd BinOpAdd BinOpAdd BinOpAdd BinOpAdd BinOpAdd Str NameLoad Str NameLoad Str NameLoad Str BinOpAdd BinOpAdd BinOpAdd NameLoad Str NameLoad Str BinOpAdd BinOpAdd BinOpAdd NameLoad Str NameLoad Str Str Str Call NameLoad NameLoad NameLoad NameLoad NameLoad Str BinOpAdd BinOpAdd Str NameLoad Str Str BinOpAdd BinOpAdd BinOpAdd BinOpAdd BinOpAdd BinOpAdd Str NameLoad Str NameLoad Str NameLoad Str BinOpAdd BinOpAdd BinOpAdd NameLoad Str NameLoad Str Str Str decorator_list Call NameLoad BinOpAdd BinOpAdd TupleLoad Num AttributeLoad NameLoad attr AttributeLoad NameLoad attr FunctionDef arguments args NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam defaults Str Str body Expr Str If NameLoad body Assign NameStore Call NameLoad BinOpMod Str Call NameLoad Assign NameStore Call NameLoad Str Call NameLoad BinOpMod Str Call NameLoad orelse Assign NameStore Call NameLoad BinOpMod Str Call NameLoad Assign NameStore Call NameLoad Str Call NameLoad BinOpMod Str Call NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Assign NameStore Call NameLoad BinOpMod Str TupleLoad NameLoad NameLoad BinOpMod Str TupleLoad NameLoad NameLoad Return BinOpMod Str Call NameLoad decorator_list Call NameLoad TupleLoad Num FunctionDef arguments args NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam defaults Str Str body Return Call NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad Lambda arguments args NameParam NameParam defaults BinOpMod Str TupleLoad NameLoad NameLoad Lambda arguments args NameParam defaults NameLoad NameLoad NameLoad decorator_list Call NameLoad AttributeLoad NameLoad attr FunctionDef arguments args NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam NameParam defaults Str Str body Expr Str Assign NameStore ListLoad Call NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad Str BinOpAdd BinOpAdd Str NameLoad Str Str Call NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad NameLoad Lambda arguments args NameParam NameParam defaults BinOpMod Str TupleLoad NameLoad NameLoad Lambda arguments args NameParam defaults BinOpMod Str NameLoad NameLoad NameLoad Str BinOpAdd BinOpAdd Str NameLoad Str Str Str If NameLoad body AugAssignAdd NameStore ListLoad BinOpMod Str Call NameLoad orelse AugAssignAdd NameStore ListLoad BinOpMod Str Call NameLoad AugAssignAdd NameStore ListLoad Str Str Return NameLoad decorator_list Call NameLoad BinOpAdd BinOpAdd TupleLoad Num AttributeLoad NameLoad attr AttributeLoad NameLoad attr
N/A N/A 
Helper routines for generating gpu kernels for nvcc.

 nvcc_kernel N/A N/A name params body N/A N/A N/A 
    Return the c code of a kernel function.

    Parameters
    ----------
    params
        The parameters to the function as one or more strings.
    body
        The [nested] list of statements for the body of the
        function. These will be separated by ';' characters.

     N/A paramstr N/A N/A ,  join params flatbody N/A N/A N/A N/A N/A b body N/A N/A N/A isinstance b N/A list tuple N/A N/A bb b N/A N/A N/A bb N/A N/A N/A b N/A N/A bodystr N/A N/A ;
 join N/A flatbody N/A N/A __global__ void %(name)s (%(paramstr)s)
    {
        %(bodystr)s;
    }
     N/A locals N/A code_version N/A N/A version N/A N/A N/A 
    Decorator to support version-based cache mechanism.

     N/A N/A N/A isinstance version tuple N/A N/A N/A TypeError version must be tuple version deco N/A N/A f N/A N/A N/A N/A f code_version version N/A f N/A N/A deco N/A N/A UNVERSIONED N/A inline_reduce N/A N/A N buf pos count manner_fn N/A N/A N/A 
    Return C++ code for a function that reduces a contiguous buffer.

    Parameters
    ----------
    N
        Length of the buffer.
    buf
        Buffer pointer.
    pos
        Index of executing thread.
    count
        Number of executing threads.
    manner_fn
        A function that accepts strings of arguments a
        and b, and returns c code for their reduction. (Example:
        return "%(a)s + %(b)s" for a sum reduction).

    :postcondition:
    This function leaves the answer in position 0 of the buffer. The
    rest of the buffer is trashed by this function.

    Notes
    ----- 
    buf should be in gpu shared memory, we access it many times.

     N/A loop_line N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[i] buf N/A r_16 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+16] N/A buf pos N/A r_8 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+8] N/A buf pos N/A r_4 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+4] N/A buf pos N/A r_2 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+2] N/A buf pos N/A r_1 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+1] N/A buf pos N/A N/A 
    {
        // This function trashes buf[1..warpSize],
        // leaving the reduction result in buf[0].

        if (%(pos)s < warpSize)
        {
            for (int i = %(pos)s + warpSize; i < %(N)s; i += warpSize)
            {
                %(buf)s[%(pos)s] = %(loop_line)s;
            }
            if (%(pos)s < 16)
            {
                //reduce so that %(pos)s 0 has the sum of everything
                if(%(pos)s + 16 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_16)s;
                if(%(pos)s + 8 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_8)s;
                if(%(pos)s + 4 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_4)s;
                if(%(pos)s + 2 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_2)s;
                if(%(pos)s + 1 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_1)s;
            }
        }
    }
     N/A locals N/A N/A code_version N/A 1 inline_reduce_max N/A N/A N buf pos count N/A N/A N/A N/A inline_reduce N buf pos count N/A N/A N/A a b N/A N/A max(%s, %s) N/A a b N/A N/A code_version N/A inline_reduce code_version inline_reduce_sum N/A N/A N buf pos count N/A N/A N/A N/A inline_reduce N buf pos count N/A N/A N/A a b N/A N/A %s + %s N/A a b N/A N/A code_version N/A inline_reduce code_version inline_reduce_min N/A N/A N buf pos count N/A N/A N/A N/A inline_reduce N buf pos count N/A N/A N/A a b N/A N/A min(%s, %s) N/A a b N/A N/A code_version N/A inline_reduce code_version inline_reduce_prod N/A N/A N buf pos count N/A N/A N/A N/A inline_reduce N buf pos count N/A N/A N/A a b N/A N/A %s * %s N/A a b N/A N/A code_version N/A inline_reduce code_version inline_softmax N/A N/A N buf buf2 threadPos threadCount N/A N/A N/A 

    Parameters
    ----------
    N
        Length of the buffer.
    threadPos
        Index of executing thread.
    threadCount
        Number of executing threads.

    :Precondition: buf and buf2 contain two identical copies of the input
        to softmax
    :Postcondition: buf contains the softmax, buf2 contains un-normalized
        softmax

    Notes
    -----
    buf and buf2 should be in gpu shared memory, we access it many times.

    We use __i as an int variable in a loop.

     N/A N/A N/A inline_reduce_max N buf threadPos threadCount __syncthreads() N/A N/A float row_max =  buf [0] __syncthreads() N/A N/A N/A N/A N/A N/A for(int __i= threadPos ; __i< N ; __i+= threadCount ){ N/A N/A N/A buf [__i] = exp( buf2 [__i] - row_max) N/A N/A N/A buf2 [__i] =  buf [__i] } __syncthreads() N/A inline_reduce_sum N buf threadPos threadCount __syncthreads() N/A N/A float row_sum =  buf [0] __syncthreads() N/A N/A N/A N/A N/A N/A for(int __i= threadPos ; __i< N ; __i+= threadCount ){ N/A N/A N/A buf [__i] =  buf2 [__i] / row_sum } __syncthreads() N/A N/A code_version N/A N/A N/A 2 N/A inline_reduce_max code_version N/A inline_reduce_sum code_version inline_reduce_fixed_shared N/A N/A N buf x stride_x pos count manner_fn manner_init b stride_b N/A N/A N/A 
    Return C++ code for a function that reduces a contiguous buffer.

    Parameters
    ----------
    N
        Length of the buffer.
    buf
        Buffer pointer of size warpSize * sizeof(float).
    pos
        Index of executing thread.
    count
        Number of executing threads.
    b
        Optional, pointer to the bias.
    stride_b
        Optional, the stride of b if b is provided.
    manner_fn
        A function that accepts strings of arguments a
        and b, and returns c code for their reduction. (Example:
        return "%(a)s + %(b)s" for a sum reduction).
    manner_init
        A function that accepts strings of arguments a
        and return c code for its initialization.

    :postcondition:
    This function leaves the answer in position 0 of the buffer. The
    rest of the buffer is trashed by this function.

    Notes
    -----
    buf should be in gpu shared memory, we access it many times.

     N/A b N/A N/A init N/A manner_init N/A %(x)s[%(pos)s * %(stride_x)s] + %(b)s[%(pos)s * %(stride_b)s] N/A locals N/A loop_line N/A manner_fn red N/A manner_init N/A %(x)s[i * %(stride_x)s] + %(b)s[i * %(stride_b)s] N/A locals N/A N/A init N/A manner_init N/A %(x)s[%(pos)s * %(stride_x)s] N/A locals N/A loop_line N/A manner_fn red N/A manner_init N/A %(x)s[i * %(stride_x)s] N/A locals N/A loop_line2 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[i] buf N/A r_16 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+16] N/A buf pos N/A r_8 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+8] N/A buf pos N/A r_4 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+4] N/A buf pos N/A r_2 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+2] N/A buf pos N/A r_1 N/A manner_fn N/A %s[%s] N/A buf pos N/A %s[%s+1] N/A buf pos N/A N/A 
    {
        // This function trashes buf[1..n_threads],
        // leaving the reduction result in buf[0].
        float red = %(init)s;
        #pragma unroll 16
        for (int i = %(pos)s + %(count)s; i<%(N)s; i += %(count)s){
          red = %(loop_line)s;
        }
        buf[%(pos)s] = red;
        __syncthreads();
        if (%(pos)s < warpSize)
        {
            for (int i = %(pos)s + warpSize; i < %(count)s; i += warpSize)
            {
                %(buf)s[%(pos)s] = %(loop_line2)s;
            }
            if (%(pos)s < 16)
            {
                //reduce so that %(pos)s 0 has the reduction of everything
                if(%(pos)s + 16 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_16)s;
                if(%(pos)s + 8 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_8)s;
                if(%(pos)s + 4 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_4)s;
                if(%(pos)s + 2 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_2)s;
                if(%(pos)s + 1 < %(N)s)
                    %(buf)s[%(pos)s] = %(r_1)s;
            }
        }
    }
     N/A locals N/A N/A code_version N/A 1 inline_reduce_fixed_shared_max N/A N/A N buf x stride_x pos count b stride_b N/A N/A N/A N/A inline_reduce_fixed_shared N buf x stride_x pos count N/A N/A N/A a b N/A N/A max(%s, %s) N/A a b N/A N/A N/A a N/A a b stride_b N/A N/A code_version N/A inline_reduce_fixed_shared code_version inline_softmax_fixed_shared N/A N/A N buf x stride_x sm sm_stride threadPos threadCount b stride_b N/A N/A N/A 

    Parameters
    ----------
    N
        Length of the buffer, atleast waprSize(32).
    buf
        A shared memory buffer of size warpSize * sizeof(float).
    x
        A ptr to the gpu memory where the row is stored.
    stride_x
        The stride between each element in x.
    sm
        A ptr to the gpu memory to store the result.
    sm_stride
        The stride between each sm element.
    threadPos
        Index of executing thread.
    threadCount
        Number of executing threads.
    b
        Optional, pointer to the bias.
    stride_b
        Optional, the stride of b if b is provided.

    :Precondition: buf is empty
    :Postcondition: buf[0] contains the softmax,
        buf2 contains un-normalized softmax

    Notes
    -----
    buf should be in gpu shared memory, we access it many times.

    We use tx as an int variable in a loop.

     N/A ret N/A N/A inline_reduce_fixed_shared_max N buf x stride_x threadPos threadCount b stride_b __syncthreads() N/A N/A float row_max =  buf [0] __syncthreads() N/A inline_reduce_fixed_shared N buf x stride_x threadPos threadCount N/A N/A N/A a b N/A N/A %s + %s N/A a b N/A N/A N/A a N/A N/A exp(%s - row_max) a b stride_b __syncthreads() N/A N/A float row_sum =  buf [0] __syncthreads() for (int tx = threadIdx.x; tx< N; tx += blockDim.x){ N/A b N/A N/A ret N/A N/A %(sm)s[tx * %(sm_stride)s] =   exp(%(x)s[tx * %(stride_x)s] +            %(b)s[tx * %(stride_b)s] - row_max) / row_sum N/A locals N/A N/A ret N/A N/A %(sm)s[tx * %(sm_stride)s] = exp(%(x)s[tx * %(stride_x)s] - row_max) / row_sum N/A locals N/A ret N/A } __syncthreads() N/A ret N/A N/A code_version N/A N/A N/A 1 N/A inline_reduce_max code_version N/A inline_reduce_sum code_version