     THCudaTensor *output = THCudaTensor_newWithSize1d(state, input->size[0]);
     dim3 blocks(input->size[0]);
     dim3 threads(MULTIMARGIN_THREADS);
-    cunn_MultiMarginCriterion_updateOutput_kernel<<<blocks,threads>>>(THCudaTensor_data(state, output),
+    if(p==1)
+      cunn_MultiMarginCriterion_updateOutput_kernel<1> <<<blocks,threads>>>(THCudaTensor_data(state, output),
+                                                                      THCudaTensor_data(state, input),
+                                                                      THCudaTensor_data(state, target),
+                                                                      input->size[0], input->size[1],
+                                                                      sizeaverage);
+    else if(p==2)
+      cunn_MultiMarginCriterion_updateOutput_kernel<2> <<<blocks,threads>>>(THCudaTensor_data(state, output),
                                                                       THCudaTensor_data(state, input),
                                                                       THCudaTensor_data(state, target),
                                                                       input->size[0], input->size[1],

      continue;

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


  }
}

template <int P>
__global__ void cunn_MultiMarginCriterion_updateGradInput_kernel(float *gradInput, float *input, float *target, int nframe, int dim, int sizeaverage)
{
  __shared__ float buffer[MULTIMARGIN_THREADS];
