   int rootId = luaT_getfieldcheckint(L, 1, "rootId") - 1;
   int maxFamilyPath = (int)luaT_getfieldcheckint(L, 1, "maxFamilyPath");
   int maxDept = (int)luaT_getfieldcheckint(L, 1, "maxDept");
+  int maxnorm = (int)luaT_getfieldcheckint(L, 1, "maxNorm");
   
   THCudaTensor *childParent = (THCudaTensor*)luaT_getfieldcheckudata(L, 1, "childParentCuda", "torch.CudaTensor");
   THCudaTensor *parentChildren = (THCudaTensor*)luaT_getfieldcheckudata(L, 1, "parentChildrenCuda", "torch.CudaTensor");

      lua_settable(L, -3);
    }
  }
    
  return 0;
}

__global__ void cunnx_SoftMaxTree_updateParameters_kernel(
  float *weight, float *bias, float *gradWeight, float *gradBias, 
  float *childParent, float *parentChildren, float *paramUpdateCuda,
  int nInput, float lr)
{
  int tx = threadIdx.x;
  int i_step = blockDim.x;
  int nodeId = paramUpdateCuda[blockIdx.x] - 1;

    THCudaTensor_data(weight), THCudaTensor_data(bias), 
    THCudaTensor_data(gradWeight), THCudaTensor_data(gradBias), 
    THCudaTensor_data(childParent), THCudaTensor_data(parentChildren), 
    THCudaTensor_data(paramUpdateCuda), weight->size[1], lr, maxnorm
  );
  
  cudaError errcode = cudaGetLastError();
  if(errcode != cudaSuccess)
    THError(cudaGetErrorString(errcode));
  
  return 0;
}

