diff options
| author | ngimel <ngimelshein@nvidia.com> | 2018-04-25 17:59:49 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2018-04-25 17:59:49 -0700 |
| commit | d48d3ef6bc4e8370a49a565ec674089acb32cf24 (patch) | |
| tree | fa34aaecebd667549563e6d3401c76798c1a1f55 /aten/src | |
| parent | 5209213fa7be89b5af2f930889d6cd595b80efad (diff) | |
| download | pytorch-d48d3ef6bc4e8370a49a565ec674089acb32cf24.tar.gz pytorch-d48d3ef6bc4e8370a49a565ec674089acb32cf24.tar.bz2 pytorch-d48d3ef6bc4e8370a49a565ec674089acb32cf24.zip | |
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
Diffstat (limited to 'aten/src')
| -rw-r--r-- | aten/src/THCUNN/SpatialGridSamplerBilinear.cu | 8 | ||||
| -rw-r--r-- | aten/src/THCUNN/VolumetricGridSamplerBilinear.cu | 12 |
2 files changed, 10 insertions, 10 deletions
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<float,Dtype>::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert<float,Dtype>::to(((iy + 1) / 2) * (IH-1)); + ix = ScalarConvert<float,Dtype>::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert<float,Dtype>::to(((iy + 1.f) / 2) * (IH-1)); // get NE, NW, SE, SW pixel values from (x, y) int ix_nw = floor(ScalarConvert<Dtype,float>::to(ix)); @@ -134,8 +134,8 @@ __global__ void SpatialGridSamplerBilinear_updateGradInput_kernel( Dtype giy = ScalarConvert<int,Dtype>::to(0); // normalize ix, iy from [-1, 1] to [0, H-1] & [0, W-1] - ix = ScalarConvert<float,Dtype>::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert<float,Dtype>::to(((iy + 1) / 2) * (IH-1));; + ix = ScalarConvert<float,Dtype>::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert<float,Dtype>::to(((iy + 1.f) / 2) * (IH-1));; // get NE, NW, SE, SW pixel values from (x, y) int ix_nw = floor(ScalarConvert<Dtype,float>::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<float,Dtype>::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert<float,Dtype>::to(((iy + 1) / 2) * (IH-1)); - iz = ScalarConvert<float,Dtype>::to(((iz + 1) / 2) * (ID-1)); + ix = ScalarConvert<float,Dtype>::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert<float,Dtype>::to(((iy + 1.f) / 2) * (IH-1)); + iz = ScalarConvert<float,Dtype>::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<int,Dtype>::to(0); // normalize ix, iy, iz from [-1, 1] to [0, IW-1] & [0, IH-1] & [0, ID-1] - ix = ScalarConvert<float,Dtype>::to(((ix + 1) / 2) * (IW-1)); - iy = ScalarConvert<float,Dtype>::to(((iy + 1) / 2) * (IH-1)); - iz = ScalarConvert<float,Dtype>::to(((iz + 1) / 2) * (ID-1)); + ix = ScalarConvert<float,Dtype>::to(((ix + 1.f) / 2) * (IW-1)); + iy = ScalarConvert<float,Dtype>::to(((iy + 1.f) / 2) * (IH-1)); + iz = ScalarConvert<float,Dtype>::to(((iz + 1.f) / 2) * (ID-1)); // get corner pixel values from (x, y, z) // for 4d, we used north-east-south-west |