     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
+    THCudaTensor_data(paramUpdateCuda), weight->size[1], lr, maxnorm
   );
   
   cudaError errcode = cudaGetLastError();
   if(errcode != cudaSuccess)
     THError(cudaGetErrorString(errcode));
-    
+  
   return 0;
 }
 

#define MINUS_LOG_THRESHOLD -18.42
#define SOFTMAXTREE_THREADS 32
#define SOFTMAXTREE_MAXCHILDREN 10000

struct addvalue_functor
{
  const float value;

  addvalue_functor(float value_) : value(value_) {}

    __host__ __device__ float operator()(const float& x) const
  {
    return (x+value);
  }
};

__global__ void cunnx_SoftMaxTree_updateOutput_kernel(
  float *output, float *logsoftOutput, float *input, float *weight, 

  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");
