 
   AT_DISPATCH_FLOATING_TYPES_AND_HALF(
       logits.scalar_type(), "SigmoidFocalLoss_backward", [&] {
-        SigmoidFocalLossBackward<scalar_t><<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-            d_logits_size, logits.contiguous().data_ptr<scalar_t>(),
-            targets.contiguous().data_ptr<int64_t>(),
-            d_losses.contiguous().data_ptr<scalar_t>(), num_classes, gamma, alpha,
-            num_samples, d_logits.data_ptr<scalar_t>());
+        SigmoidFocalLossBackward<scalar_t>
+            <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
+                d_logits_size, logits.contiguous().data_ptr<scalar_t>(),
+                targets.contiguous().data_ptr<int64_t>(),
+                d_losses.contiguous().data_ptr<scalar_t>(), num_classes, gamma,
+                alpha, num_samples, d_logits.data_ptr<scalar_t>());
       });
 
   THCudaCheck(cudaGetLastError());

  AT_ASSERTM(logits.device().is_cuda(), "logits must be a CUDA tensor");
  AT_ASSERTM(targets.device().is_cuda(), "targets must be a CUDA tensor");
  AT_ASSERTM(logits.dim() == 2, "logits should be NxClass");

  const int num_samples = logits.size(0);



  AT_DISPATCH_FLOATING_TYPES_AND_HALF(
      logits.scalar_type(), "SigmoidFocalLoss_forward", [&] {
        SigmoidFocalLossForward<scalar_t>
            <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
                losses_size, logits.contiguous().data_ptr<scalar_t>(),
                targets.contiguous().data_ptr<int64_t>(), num_classes, gamma,
                alpha, num_samples, losses.data_ptr<scalar_t>());
      });
  THCudaCheck(cudaGetLastError());
  return losses;
