 
 #define MULTIMARGIN_THREADS 128
 
+template <int P>
 __global__ void cunn_MultiMarginCriterion_updateOutput_kernel(float *output, float *input, float *target, int nframe, int dim, int sizeaverage)
 {
   __shared__ float buffer[MULTIMARGIN_THREADS];


    THCudaStorage_fill(state, target, target_);

    cunn_MultiMarginCriterion_updateOutput_kernel<<<blocks,threads>>>(output->data,
                                                                      THCudaTensor_data(state, input),
                                                                      target->data,
                                                                      1, input->size[0],


    if(z > 0)
    {
      float h = (P == 1) ? g : 2*g*z;
      buffer[threadIdx.x] -= h;
      gradInput_k[i] = h;
    }
    else
      gradInput_k[i] = 0;
