From 15f097e69969b8437ed8f81fb6ba228d15db75f5 Mon Sep 17 00:00:00 2001
From: Francisco Massa <fvsmassa@gmail.com>
Date: Thu, 17 Oct 2019 10:02:20 +0200
Subject: [PATCH] Fix CUDA builds on Windows

---
 torchvision/csrc/cuda/PSROIAlign_cuda.cu | 10 ++++++++--
 torchvision/csrc/cuda/PSROIPool_cuda.cu  | 10 ++++++++--
 2 files changed, 16 insertions(+), 4 deletions(-)

diff --git a/torchvision/csrc/cuda/PSROIAlign_cuda.cu b/torchvision/csrc/cuda/PSROIAlign_cuda.cu
index 3237c7f7689..5bf86eaedb9 100644
--- a/torchvision/csrc/cuda/PSROIAlign_cuda.cu
+++ b/torchvision/csrc/cuda/PSROIAlign_cuda.cu
@@ -336,7 +336,10 @@ std::tuple<at::Tensor, at::Tensor> PSROIAlign_forward_cuda(
 
   cudaStream_t stream = at::cuda::getCurrentCUDAStream();
 
-  dim3 grid(std::min(at::cuda::ATenCeilDiv(output_size, 512L), 4096L));
+  dim3 grid(std::min(
+      at::cuda::ATenCeilDiv(
+          static_cast<int64_t>(output_size), static_cast<int64_t>(512)),
+      static_cast<int64_t>(4096)));
   dim3 block(512);
 
   AT_DISPATCH_FLOATING_TYPES_AND_HALF(
@@ -395,7 +398,10 @@ at::Tensor PSROIAlign_backward_cuda(
 
   cudaStream_t stream = at::cuda::getCurrentCUDAStream();
 
-  dim3 grid(std::min(at::cuda::ATenCeilDiv(grad.numel(), 512L), 4096L));
+  dim3 grid(std::min(
+      at::cuda::ATenCeilDiv(
+          static_cast<int64_t>(grad.numel()), static_cast<int64_t>(512)),
+      static_cast<int64_t>(4096)));
   dim3 block(512);
 
   // handle possibly empty gradients
diff --git a/torchvision/csrc/cuda/PSROIPool_cuda.cu b/torchvision/csrc/cuda/PSROIPool_cuda.cu
index a2c5addd936..fabbe5bf63e 100644
--- a/torchvision/csrc/cuda/PSROIPool_cuda.cu
+++ b/torchvision/csrc/cuda/PSROIPool_cuda.cu
@@ -173,7 +173,10 @@ std::tuple<at::Tensor, at::Tensor> PSROIPool_forward_cuda(
 
   cudaStream_t stream = at::cuda::getCurrentCUDAStream();
 
-  dim3 grid(std::min(at::cuda::ATenCeilDiv(output_size, 512L), 4096L));
+  dim3 grid(std::min(
+      at::cuda::ATenCeilDiv(
+          static_cast<int64_t>(output_size), static_cast<int64_t>(512)),
+      static_cast<int64_t>(4096)));
   dim3 block(512);
 
   AT_DISPATCH_FLOATING_TYPES_AND_HALF(
@@ -229,7 +232,10 @@ at::Tensor PSROIPool_backward_cuda(
 
   cudaStream_t stream = at::cuda::getCurrentCUDAStream();
 
-  dim3 grid(std::min(at::cuda::ATenCeilDiv(grad.numel(), 512L), 4096L));
+  dim3 grid(std::min(
+      at::cuda::ATenCeilDiv(
+          static_cast<int64_t>(grad.numel()), static_cast<int64_t>(512)),
+      static_cast<int64_t>(4096)));
   dim3 block(512);
 
   // handle possibly empty gradients