   cunnx_SoftMaxTree_accGradParameters_kernel<<<blocks,threads>>>(
     THCudaTensor_data(gradWeight), THCudaTensor_data(Bias), 
     THCudaTensor_data(input), THCudaTensor_data(linearGradOutput), 
-    THCudaTensor_data(target), THCudaTensor_data(childParent), 
-    THCudaTensor_data(parentChildren), 
-    input->size[1], rootId, maxFamilyPath, scale
+    THCudaTensor_data(nodeUpdateCuda), THCudaTensor_data(target), 
+    THCudaTensor_data(childParent), THCudaTensor_data(parentChildren), 
+    input->size[1], rootId, maxFamilyPath, maxDept, scale 
   );
+  
+  // copy nodeIds from device to host
+  
+  /* updates will contain parentId (key) sum of scales (value)*/
+  lua_pushinteger(L, (int)(parentId+1));
+  lua_gettable(L, -2);
+  double count = lua_tonumber(L, -1) + scale;
+  lua_pop(L, 1);
+  
+  lua_pushinteger(L, (int)(parentId+1)); /* key */
+  lua_pushnumber(L, count); /* value */
+  lua_settable(L, -3);
     
   return 0;
 }

      nodeGradBias[i] += scale*nodeGrad[i]
    }
    
    n += nChildren;
    /* Break when root is reached */
    if (parentId == rootId)
    {
      break;
    }
    childId = parentId;



__global__ void cunnx_SoftMaxTree_updateOutput_kernel(
  float *output, float *logsoftOutput, float *input, float *weight, 
  float *bias, float *target, float *childParent, float *parentChildren, 
  int nInput, int rootId, int maxFamilyPath)
{
  //__shared__ float input_buffer[nInput]; // constant might be faster
