 cunn_LogSoftMax_updateOutput_kernel(float *output, float *input, int classes)
 {
   extern __shared__ float buffer[];
+  // forward pointers to batch[blockIdx.x]
+  // each block handles a sample in the mini-batch
   input += blockIdx.x * classes;
   output += blockIdx.x * classes;
 
+  // find the max of the batch
   float threadMax =
     ilpReduce<MaxFloat, ILP>(input, classes, MaxFloat(), -FLT_MAX);
+  // find the max over all batches
   float max_k =
     blockReduce<MaxFloat>(buffer, threadMax, MaxFloat(), -FLT_MAX);
 

  float warpVal = defaultVal;

  // First warp will perform per-warp reductions for the remaining warps
  if ((threadIdx.x / 32) == 0)
  {
    int lane = threadIdx.x % 32;

    if (lane < blockDim.x / 32)
    {
#pragma unroll
      for (int i = 0; i < 32; ++i)

            float defaultVal,
            const Finalize& f)
{
  // To avoid RaW races from chaining blockReduce calls together, we
  // need a sync here
  __syncthreads();

