From d48d3ef6bc4e8370a49a565ec674089acb32cf24 Mon Sep 17 00:00:00 2001 From: ngimel Date: Wed, 25 Apr 2018 17:59:49 -0700 Subject: Make cuda 9 behave as cuda 8 wrt half conversions (#6958) * Make cuda 9 behave as cuda 8 wrt half conversions Cuda 9 is too smart about implicit half conversions, this would disable them so that cuda 8 and cuda 9 behave in the same way wrt half. * try fixing windows build * one more broken conversion --- aten/src/THCUNN/SpatialGridSamplerBilinear.cu | 8 ++++---- aten/src/THCUNN/VolumetricGridSamplerBilinear.cu | 12 ++++++------ 2 files changed, 10 insertions(+), 10 deletions(-) (limited to 'aten/src') diff --git a/aten/src/THCUNN/SpatialGridSamplerBilinear.cu b/aten/src/THCUNN/SpatialGridSamplerBilinear.cu index 3bbc10557f..30a1a5d5ad 100644 --- a/aten/src/THCUNN/SpatialGridSamplerBilinear.cu +++ b/aten/src/THCUNN/SpatialGridSamplerBilinear.cu @@ -52,8 +52,8 @@ __global__ void SpatialGridSamplerBilinear_updateOutput_kernel( Dtype iy = grid[n][h][w][1]; // normalize ix, iy from [-1, 1] to [0, IH-1] & [0, IW-1] - ix = ScalarConvert::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert::to(((iy + 1) / 2) * (IH-1)); + ix = ScalarConvert::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert::to(((iy + 1.f) / 2) * (IH-1)); // get NE, NW, SE, SW pixel values from (x, y) int ix_nw = floor(ScalarConvert::to(ix)); @@ -134,8 +134,8 @@ __global__ void SpatialGridSamplerBilinear_updateGradInput_kernel( Dtype giy = ScalarConvert::to(0); // normalize ix, iy from [-1, 1] to [0, H-1] & [0, W-1] - ix = ScalarConvert::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert::to(((iy + 1) / 2) * (IH-1));; + ix = ScalarConvert::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert::to(((iy + 1.f) / 2) * (IH-1));; // get NE, NW, SE, SW pixel values from (x, y) int ix_nw = floor(ScalarConvert::to(ix)); diff --git a/aten/src/THCUNN/VolumetricGridSamplerBilinear.cu b/aten/src/THCUNN/VolumetricGridSamplerBilinear.cu index 6fa8ae1fd3..43b8ceff1c 100644 --- a/aten/src/THCUNN/VolumetricGridSamplerBilinear.cu +++ b/aten/src/THCUNN/VolumetricGridSamplerBilinear.cu @@ -56,9 +56,9 @@ __global__ void VolumetricGridSamplerBilinear_updateOutput_kernel( Dtype iz = grid[n][d][h][w][2]; // normalize ix, iy, iz from [-1, 1] to [0, IW-1] & [0, IH-1] & [0, ID-1] - ix = ScalarConvert::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert::to(((iy + 1) / 2) * (IH-1)); - iz = ScalarConvert::to(((iz + 1) / 2) * (ID-1)); + ix = ScalarConvert::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert::to(((iy + 1.f) / 2) * (IH-1)); + iz = ScalarConvert::to(((iz + 1.f) / 2) * (ID-1)); // get corner pixel values from (x, y, z) // for 4d, we used north-east-south-west @@ -201,9 +201,9 @@ __global__ void VolumetricGridSamplerBilinear_updateGradInput_kernel( Dtype giz = ScalarConvert::to(0); // normalize ix, iy, iz from [-1, 1] to [0, IW-1] & [0, IH-1] & [0, ID-1] - ix = ScalarConvert::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert::to(((iy + 1) / 2) * (IH-1)); - iz = ScalarConvert::to(((iz + 1) / 2) * (ID-1)); + ix = ScalarConvert::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert::to(((iy + 1.f) / 2) * (IH-1)); + iz = ScalarConvert::to(((iz + 1.f) / 2) * (ID-1)); // get corner pixel values from (x, y, z) // for 4d, we used north-east-south-west -- cgit v1.2.3