code: def inline_softmax(
      N, buf, buf2, threadPos, threadCount, dtype="float32"
  ):
      """
      Generate code for a softmax.
  
      On entry, `buf` and `buf2` must contain two identical copies of
      the input to softmax.
  
      After the code returns `buf` contains the softmax, `buf2` contains
      un-normalized softmax.
  
      Parameters
      ----------
      N
          Length of the buffer.
      threadPos
          Index of executing thread.
      threadCount
          Number of executing threads.
      dtype
          Dtype of the softmax's output.
  
      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.
  
      """
      ctype = gpuarray.dtype_to_ctype(dtype)
      # get max of buf (trashing all but buf[0])
      return [
          inline_reduce_max(N, buf, threadPos, threadCount),
          "__syncthreads()",
          ("%s row_max = " + buf + "[0]") % ctype,
          "__syncthreads()",
          "for(int __i="
          + threadPos
          + "; __i<"
          + N
          + "; __i+="
          + threadCount
          + "){",
          buf + "[__i] = exp(" + buf2 + "[__i] - row_max)",
          buf2 + "[__i] = " + buf + "[__i]",
          "}",
          "__syncthreads()",
          inline_reduce_sum(N, buf, threadPos, threadCount),
          "__syncthreads()",
          ("%s row_sum = " + buf + "[0]") % ctype,
          "__syncthreads()",
          # divide each exp() result by the sum to complete the job.
          "for(int __i="
          + threadPos
          + "; __i<"
          + N
          + "; __i+="
          + threadCount
          + "){",
          buf + "[__i] = " + buf2 + "[__i] / row_sum",
          "}",
          "__syncthreads()",
      ]
  
file_path: Theano/theano/gpuarray/kernel_codegen.py
imports: from __future__ import absolute_import, print_function, division
label_desc: Use image node_id 8 for calling a global function with example usage: inline_softmax(N, buf, buf2, threadPos, threadCount, dtype) and returns: list