@code_version((1,) + inline_reduce_max.code_version +
              inline_reduce_sum.code_version)
def inline_softmax_fixed_shared(N, buf, x, stride_x,
                                sm, sm_stride,
                                threadPos, threadCount,
                                b='', stride_b=''):
    """

    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.

    """
    ret = [
        # get max of buf (trashing all but buf[0])
        inline_reduce_fixed_shared_max(N, buf, x, stride_x,
                                       threadPos, threadCount, b, stride_b),
        '__syncthreads()',
        'float row_max = ' + buf + '[0]',
        '__syncthreads()',
        inline_reduce_fixed_shared(N, buf, x, stride_x, threadPos, threadCount,
                                   lambda a, b: "%s + %s" % (a, b),
                                   lambda a: "exp(%s - row_max)" % a,
                                   b, stride_b),
        '__syncthreads()',
        'float row_sum = ' + buf + '[0]',
        '__syncthreads()',
        "for (int tx = threadIdx.x; tx< N; tx += blockDim.x){",
        ]
    # This set all value correctly
    if b:
        ret += [
            "%(sm)s[tx * %(sm_stride)s] = "
            "  exp(%(x)s[tx * %(stride_x)s] +"
            "            %(b)s[tx * %(stride_b)s] - row_max)"
            " / row_sum" % locals()]
    else:
        ret += [
            "%(sm)s[tx * %(sm_stride)s] = "
            "exp(%(x)s[tx * %(stride_x)s] - row_max) / row_sum" % locals()]
    ret += [
        "}",
        '__syncthreads()',
    ]
    return ret
dataset/ETHPy150Open rizar/attention-lvcsr/libs/Theano/theano/sandbox/cuda/kernel_codegen.py/inline_softmax_fixed_shared/original