   buffer[threadIdx.x] = -FLT_MAX;
   for (int i=i_start; i<i_end; i+=i_step)
   {
-    float z = input_k[i];
+    float z = input_k[i*stride];
     if(buffer[threadIdx.x] < z)
       buffer[threadIdx.x] = z;
   }

  // sum?
  buffer[threadIdx.x] = 0;
  for (int i=i_start; i<i_end; i+=i_step)
    buffer[threadIdx.x] += gradOutput_k[i] * output_k[i];

  __syncthreads();


#define MINUS_LOG_THRESHOLD -18.42
#define SOFTMAX_THREADS 128

__global__ void cunn_SoftMax_updateOutput_kernel(float *output, float *input,
                                                 int nframe, int dim, int stride)
{
  __shared__ float buffer[SOFTMAX_THREADS+1];
  float *input_k = input + blockIdx.x*dim*stride + blockIdx.y;
  float *output_k = output + blockIdx.x*dim*stride + blockIdx.y;

  int i_start = threadIdx.x;
  int i_end = dim;
