   float max_k = buffer[SOFTMAX_THREADS];
   buffer[threadIdx.x] = 0;
   for (int i=i_start; i<i_end; i+=i_step) {
-    float z = __expf(input_k[i]-max_k);
+    float z = __expf(input_k[i*stride]-max_k);
     buffer[threadIdx.x] += z;
-    output_k[i] = z;
+    output_k[i*stride] = z;
   }
 
   __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)
{
  __shared__ float buffer[SOFTMAX_THREADS+1];
  int k = blockIdx.x;
  float *input_k = input + k*dim;
  float *output_k = output + k*dim;

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

  THCudaTensor_free(state, input);
  return 1;
}

static int cunn_SoftMax_updateGradInput(lua_State *L)
{
