   }
   else
   {
-    dim3 blocks(batchSize, stride);
-    dim3 threads(LOGSOFTMAX_THREADS);
-    cunn_SpatialLogSoftMax_updateGradInput_kernel<<<blocks,threads,
-      0, THCState_getCurrentStream(state)>>>(THCudaTensor_data(state, gradInput),
-                                             THCudaTensor_data(state, output),
-                                             THCudaTensor_data(state, gradOutput),
-                                             batchSize, classSize, stride);
+    dim3 grid(batchSize);
+    dim3 block(1024);
+
+    cunn_SpatialLogSoftMax_updateGradInput_kernel
+      <<<grid, block, 0, THCState_getCurrentStream(state)>>>(
+        THCudaTensor_data(state, gradInput),
+        THCudaTensor_data(state, output),
+        THCudaTensor_data(state, gradOutput),
+        classSize, height, width
+    );
   }
 
   cudaError errcode = cudaGetLastError();

  }
  else
  {
    dim3 blocks(batchSize, stride);
    dim3 threads(LOGSOFTMAX_THREADS);
    cunn_SpatialLogSoftMax_updateOutput_kernel<<<blocks,threads,
      0, THCState_getCurrentStream(state)>>>(THCudaTensor_data(state, output),
                                             THCudaTensor_data(state, input),
                                             batchSize, classSize, stride);
  }
  cudaError errcode = cudaGetLastError();
  if (errcode != cudaSuccess)
  {

  THCudaTensor_free(state, gradOutput);
  THCudaTensor_free(state, output);
}
