   
   luaL_argcheck(L, input->nDimension == 2, 2, "2D(batch mode) tensor expected");
   luaL_argcheck(L, input->size[1] == inputSize, 2, "invalid input size");  
+  luaL_argcheck(L, maxFamily <= SOFTMAXTREE_MAXCHILDREN, 2, "Hierarchy has node(s) with too many children");
   
   input = THCudaTensor_newContiguous(input);
   THCudaTensor_resize1d(output, input->size[0]);
+  
+  // for some reason, this is necessary:
+  THCudaTensor_zero(logsoftOutput);
   
   /* call cudakernel */
   dim3 blocks(input->size[0]); // each block is an example

    // AKA linearGradOutput (we reuse the _multiBuffer Tensor)
    nodeGrad = logsoftOutput + maxFamilyPath*k + n; 

    for(int i = tx; i < nChildren; i+=i_step)
      nodeGrad[i] = -exp(nodeGrad[i])*grad;
    
    __syncthreads();
    if (tx == 0) // compare this to % childIdx
      nodeGrad[childIdx] += grad;

    /* Linear */
    nodeWeight = weight + parentIdx*nInput;

  int inputSize = luaT_getfieldcheckint(L, 1, "inputSize");
  int rootId = luaT_getfieldcheckint(L, 1, "rootId") - 1;
  int maxFamilyPath = (int)luaT_getfieldcheckint(L, 1, "maxFamilyPath");
  int maxFamily = (int)luaT_getfieldcheckint(L, 1, "maxFamily");
  
  THCudaTensor *childParent = (THCudaTensor*)luaT_getfieldcheckudata(L, 1, "childParentCuda", "torch.CudaTensor");
  THCudaTensor *parentChildren = (THCudaTensor*)luaT_getfieldcheckudata(L, 1, "parentChildrenCuda", "torch.CudaTensor");
