       for (int j=0; j<nChildren; j++)
       {
         // multiply accumulate weights
-        nodeGradWeight[j*nInput + i] += scale*nodeGrad[j]*buffer[tx];
+        nodeGradWeight[j*nInput + i] += scale*nodeGradOutput[j]*buffer[tx];
       }
     }
     
     // cadd bias
-    for (int i=tx; i<nChildren; i+=i_step)
-    {
-      // multiply accumulate weights
-      nodeGradBias[i] += scale*nodeGrad[i]
+    for (int j=tx; j<nChildren; j+=i_step)
+    {
+      // multiply accumulate biases
+      nodeGradBias[j] += scale*nodeGradOutput[j];
     }
     
     // keep track of which node gets gradients

  float *childParent, float *parentChildren, 
  int nInput, int rootId, int maxFamilyPath, int maxDept, float scale)
{
  __shared__ float buffer[SOFTMAXTREE_THREADS+1];
  int tx = threadIdx.x;
  int i_step = blockDim.x;
  int k = blockIdx.x;

  THCudaTensor *gradWeight = (THCudaTensor*)luaT_getfieldcheckudata(L, 1, "gradWeight", "torch.CudaTensor");
  THCudaTensor *gradBias = (THCudaTensor*)luaT_getfieldcheckudata(L, 1, "gradBias", "torch.CudaTensor");
  
  int i, j;
  THIntTensor *nodeUpdate;
  
  lua_getfield(L, 1, "updates");
  
  luaL_argcheck(L, input->nDimension == 2, 2, "2D(batch mode) tensor expected");
  luaL_argcheck(L, input->size[1] == inputSize, 2, "invalid input size");  
