   return 1;
 }
 
+__global__ void cunnx_SoftMaxTree_accGradParameters_kernel(
+  float *gradWeight, float *gradBias, float *input, float* linearGradOutput,
+  float *target, float* childParent, float* parentChildren, 
+  float scale, int nInput, int rootId, int maxFamilyPath)
+{
+  __shared__ float buffer[SOFTMAXTREE_THREADS+1];
+  int tx = threadIdx.x;
+  int i_step = blockDim.x;
+  int k = blockIdx.x;
+  float *input_k = input + k*nInput;
+  float *nodeGradOutput, *nodeGradWeight, *nodeGradBias;
+  int childId = target[k] - 1;
+  int parentId, parentIdx, childIdx, nChildren;
+  float *node;
+  int n = 0;
+  THIntTensor *node;
+  
+  // loop through nodes
+  while(1)
+  {
+    /* get next Node in Tree */
+    node = childParent + childId*2;
+    parentId = (int)node[0] - 1;
+    childIdx = (int)node[1] - 1;
+    
+    node = parentChildren + parentId*2;
+    parentIdx = (int)node[0] - 1;
+    nChildren = (int)node[1];
+    
+    nodeGradOutput = linearGradOutput + maxFamilyPath*k + n; 
+    nodeGradWeight = gradWeight + parentIdx*nInput;
+    nodeGradBias = gradBias + parentIdx;
+      
+    THTensor_(addr)(nodeGradWeight, 1, nodeGradWeight, scale, nodeGradOutput, nodeInput);
+    THTensor_(cadd)(nodeGradBias, nodeGradBias, scale, nodeGradOutput);
+    
+    // addr weights (scalar-products)
+    for (int i=tx; i<nInput; i+=i_step)
+    {
+      // copy input to buffer
+      buffer[tx] = input_k[i];
+    
+      for (int j=0; j<nChildren; j++)
+      {
+        // multiply accumulate weights
+        nodeGradWeight[j*nInput + i] += scale*nodeGrad[j]*buffer[tx];
+      }
+    }
+    
+    // cadd bias
+    for (int i=tx; i<nChildren; i+=i_step)
+    {
+      // multiply accumulate weights
+      nodeGradBias[i] += scale*nodeGrad[i]
+    }
+    
+    n += nChildren;
+    /* Break when root is reached */
+    if (parentId == rootId)
+    {
+      break;
+    }
+    childId = parentId;
+  }
+}
+
 static int cunnx_SoftMaxTree_accGradParameters(lua_State *L)
 {
   THCudaTensor *input = (THCudaTensor*)luaT_checkudata(L, 2, "torch.CudaTensor");  

  
  luaL_argcheck(L, input->nDimension == 2, 2, "2D(batch mode) tensor expected");
  luaL_argcheck(L, input->size[1] == inputSize, 2, "invalid input size");  
    
  return 0;
}

        // multiply
        buffer[tx] += nodeGrad[j]*nodeWeight[j*nInput + i];
      }
      // accumulate into global memory
      gradInput_k[i] += buffer[tx];
    }
    
    n += nChildren;
    /* Break when root is reached */
