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

  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)

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);

