code: def inline_softmax_fixed_shared(
      N,
      buf,
      x,
      stride_x,
      load_x,
      sm,
      sm_stride,
      write_sm,
      threadPos,
      threadCount,
      b="",
      stride_b="",
      load_b="",
      dtype="float32",
  ):
      """
      Generate code to perform softmax with a fixed amount of shared
      memory.
  
      On entry, `buf` is assumed to be empty.
  
      On exit, `buf[0]` contains the softmax, `buf2` contains
      un-normalized softmax.
  
      Parameters
      ----------
      N
          Length of the buffer, atleast waprSize(32).
      buf
          A shared memory buffer of size warpSize * sizeof(dtype).
      x
          A ptr to the gpu memory where the row is stored.
      stride_x
          The stride between each element in x.
      load_x
          Wrapper to read from x.
      sm
          A ptr to the gpu memory to store the result.
      sm_stride
          The stride between each sm element.
      write_sm
          Wrapper before writing to sm.
      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.
      load_b
          Optional, wrapper to read from b if b is provided.
      dtype
          Optional, the dtype of the softmax's output if not float32.
  
      Notes
      -----
      `buf` should be in gpu shared memory, we access it many times.
  
      We use tx as an int variable in a loop.
  
      """
      ctype = gpuarray.dtype_to_ctype(dtype)
      ret = [
          # get max of buf (trashing all but buf[0])
          inline_reduce_fixed_shared_max(
              N,
              buf,
              x,
              stride_x,
              load_x,
              threadPos,
              threadCount,
              b,
              stride_b,
              load_b,
              dtype,
          ),
          "__syncthreads()",
          ("%s row_max = " + buf + "[0]") % ctype,
          "__syncthreads()",
          inline_reduce_fixed_shared(
              N,
              buf,
              x,
              stride_x,
              load_x,
              threadPos,
              threadCount,
              lambda a, b: "%s + %s" % (a, b),
              lambda a: "exp(%s - row_max)" % a,
              b,
              stride_b,
              load_b,
              dtype,
          ),
          "__syncthreads()",
          ("%s row_sum = " + buf + "[0]") % ctype,
          "__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] = "
              "  %(write_sm)s(exp(%(load_x)s(%(x)s[tx * %(stride_x)s]) +"
              "            %(load_b)s(%(b)s[tx * %(stride_b)s]) - row_max)"
              " / row_sum)" % locals()
          ]
      else:
          ret += [
              "%(sm)s[tx * %(sm_stride)s] = "
              "%(write_sm)s(exp(%(load_x)s(%(x)s[tx * %(stride_x)s]) - row_max)"
              " / row_sum)" % locals()
          ]
      ret += [
          "}",
          "__syncthreads()",
      ]
      return ret
  
file_path: Theano/theano/gpuarray/kernel_codegen.py
imports: from __future__ import absolute_import, print_function, division
label_desc: Use image node_id 11 for calling a global function with example usage: inline_softmax_fixed_shared(N, buf, x, stride_x, load_x, sm, sm_stride, write_sm, threadPos, threadCount, b, stride_b, load_b, dtype) and returns: ret