 
     THCudaStorage_fill(state, target, target_);
 
-    cunn_MultiMarginCriterion_updateOutput_kernel<<<blocks,threads>>>(output->data,
+    if(p==1)
+      cunn_MultiMarginCriterion_updateOutput_kernel<1> <<<blocks,threads>>>(output->data,
+                                                                      THCudaTensor_data(state, input),
+                                                                      target->data,
+                                                                      1, input->size[0],
+                                                                      sizeaverage);
+    else if(p==2)
+      cunn_MultiMarginCriterion_updateOutput_kernel<2> <<<blocks,threads>>>(output->data,
                                                                       THCudaTensor_data(state, input),
                                                                       target->data,
                                                                       1, input->size[0],

      continue;

    if(z > 0)
      buffer[threadIdx.x] += z;
  }
  __syncthreads();


{
  THCState *state = getCutorchState(L);
  THCudaTensor *input = (THCudaTensor*)luaT_checkudata(L, 2, "torch.CudaTensor");
  int p = luaT_getfieldchecknumber(L, 1, "p");
  int sizeaverage = luaT_getfieldcheckboolean(L, 1, "sizeAverage");

  input = THCudaTensor_newContiguous(state, input);
