From 9c9816a7765bdfe6a482ae1ae796c560283ffaaa Mon Sep 17 00:00:00 2001 From: Michael Wootton Date: Fri, 21 Sep 2018 14:49:10 -0500 Subject: [PATCH 1/6] MIOpen: Batchnorm - Allow half/half and half/float, disallow double --- aten/src/ATen/native/Normalization.cpp | 3 +-- aten/src/ATen/native/miopen/BatchNorm_miopen.cpp | 4 +--- 2 files changed, 2 insertions(+), 5 deletions(-) diff --git a/aten/src/ATen/native/Normalization.cpp b/aten/src/ATen/native/Normalization.cpp index 24d8a41fb50271..b44156becc2dc3 100644 --- a/aten/src/ATen/native/Normalization.cpp +++ b/aten/src/ATen/native/Normalization.cpp @@ -63,8 +63,7 @@ Tensor batch_norm( } bool use_miopen = (input.type().is_cuda() - && (input.type().scalarType() != at::kHalf - || weight.type().scalarType() == at::kFloat) + && input.type().scalarType() != at::kDouble && weight.defined() && bias.defined() && ((running_mean.defined() && running_var.defined()) || (!running_mean.defined() && !running_var.defined() && training)) diff --git a/aten/src/ATen/native/miopen/BatchNorm_miopen.cpp b/aten/src/ATen/native/miopen/BatchNorm_miopen.cpp index 997431b7a86170..f79ca1eed20b08 100644 --- a/aten/src/ATen/native/miopen/BatchNorm_miopen.cpp +++ b/aten/src/ATen/native/miopen/BatchNorm_miopen.cpp @@ -67,9 +67,7 @@ std::tuple miopen_batch_norm( checkAllDefined(c, {running_mean, running_var}); } checkAllSameGPU(c, {input, weight, bias, running_mean, running_var}); - if (input->type().scalarType() == ScalarType::Half) { - checkScalarType(c, weight, ScalarType::Float); - } else { + if (input->type().scalarType() != ScalarType::Half) { checkAllSameType(c, {input, weight}); } checkAllSameType(c, {weight, bias, running_mean, running_var}); From 326c999487dcf92cc66980da3025e8aa74ec7e1e Mon Sep 17 00:00:00 2001 From: Michael Wootton Date: Fri, 21 Sep 2018 16:04:15 -0500 Subject: [PATCH 2/6] MIOpen: Honor DIM_MAX --- aten/src/ATen/native/Convolution.cpp | 4 +++- aten/src/ATen/native/Normalization.cpp | 4 ++++ 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/aten/src/ATen/native/Convolution.cpp b/aten/src/ATen/native/Convolution.cpp index f359d67c72e786..c9ff07861eb169 100644 --- a/aten/src/ATen/native/Convolution.cpp +++ b/aten/src/ATen/native/Convolution.cpp @@ -3,6 +3,8 @@ #include "ATen/Config.h" +#define MIOPEN_DIM_MAX = 4 + namespace at { namespace native { struct ConvParams { @@ -120,7 +122,7 @@ auto ConvParams::use_cudnn(const at::Tensor& input) const -> bool { } auto ConvParams::use_miopen(const at::Tensor& input) const -> bool { - if (!detail::getCUDAHooks().compiledWithMIOpen() || !input.type().is_cuda() || !cudnn_enabled) + if (!detail::getCUDAHooks().compiledWithMIOpen() || !input.type().is_cuda() || input.dim() > MIOPEN_DIM_MAX || getenv("DISABLE_MIOPEN") != NULL) return false; return true; } diff --git a/aten/src/ATen/native/Normalization.cpp b/aten/src/ATen/native/Normalization.cpp index b44156becc2dc3..a54333707b0d66 100644 --- a/aten/src/ATen/native/Normalization.cpp +++ b/aten/src/ATen/native/Normalization.cpp @@ -7,6 +7,8 @@ #include +#define MIOPEN_DIM_MAX 4 + namespace at { namespace native { namespace { @@ -63,11 +65,13 @@ Tensor batch_norm( } bool use_miopen = (input.type().is_cuda() + && input.dim() < MIOPEN_DIM_MAX && input.type().scalarType() != at::kDouble && weight.defined() && bias.defined() && ((running_mean.defined() && running_var.defined()) || (!running_mean.defined() && !running_var.defined() && training)) && detail::getCUDAHooks().compiledWithMIOpen() + && getenv("DISABLE_MIOPEN") == NULL ); if (use_miopen) { From 5f219395698c8519b3479ea96fe90ac6b2de3b5b Mon Sep 17 00:00:00 2001 From: Michael Wootton Date: Mon, 24 Sep 2018 02:17:03 -0500 Subject: [PATCH 3/6] Change define to a static. --- aten/src/ATen/native/Convolution.cpp | 2 +- aten/src/ATen/native/Normalization.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/aten/src/ATen/native/Convolution.cpp b/aten/src/ATen/native/Convolution.cpp index 154b08070d16cf..e52af4e164eafa 100644 --- a/aten/src/ATen/native/Convolution.cpp +++ b/aten/src/ATen/native/Convolution.cpp @@ -3,7 +3,7 @@ #include "ATen/Config.h" -#define MIOPEN_DIM_MAX = 4 +static int MIOPEN_DIM_MAX = 4; namespace at { namespace native { diff --git a/aten/src/ATen/native/Normalization.cpp b/aten/src/ATen/native/Normalization.cpp index 5112c17ef19e3b..7831428c01e3a6 100644 --- a/aten/src/ATen/native/Normalization.cpp +++ b/aten/src/ATen/native/Normalization.cpp @@ -7,7 +7,7 @@ #include -#define MIOPEN_DIM_MAX 4 +static int MIOPEN_DIM_MAX = 4; namespace at { namespace native { From 3f275fc455590d24e6ffeb08593a56ede1de2da6 Mon Sep 17 00:00:00 2001 From: Michael Wootton Date: Tue, 25 Sep 2018 13:36:18 -0500 Subject: [PATCH 4/6] Cache MIOPEN_DISABLED --- aten/src/ATen/native/Convolution.cpp | 3 ++- aten/src/ATen/native/Normalization.cpp | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/aten/src/ATen/native/Convolution.cpp b/aten/src/ATen/native/Convolution.cpp index e52af4e164eafa..f8f6dd227885ce 100644 --- a/aten/src/ATen/native/Convolution.cpp +++ b/aten/src/ATen/native/Convolution.cpp @@ -4,6 +4,7 @@ #include "ATen/Config.h" static int MIOPEN_DIM_MAX = 4; +static bool MIOPEN_ENABLED = getenv("DISABLE_MIOPEN") != NULL; namespace at { namespace native { @@ -126,7 +127,7 @@ auto ConvParams::use_miopen(const at::Tensor& input) const -> bool { && detail::getCUDAHooks().compiledWithMIOpen() && input.type().is_cuda() && input.dim() > MIOPEN_DIM_MAX - && getenv("DISABLE_MIOPEN") != NULL + && MIOPEN_ENABLED ; } diff --git a/aten/src/ATen/native/Normalization.cpp b/aten/src/ATen/native/Normalization.cpp index 7831428c01e3a6..4b39dcd46e981f 100644 --- a/aten/src/ATen/native/Normalization.cpp +++ b/aten/src/ATen/native/Normalization.cpp @@ -8,6 +8,7 @@ #include static int MIOPEN_DIM_MAX = 4; +static bool MIOPEN_ENABLED = getenv("DISABLE_MIOPEN") != NULL; namespace at { namespace native { @@ -75,7 +76,7 @@ Tensor batch_norm( && ((running_mean.defined() && running_var.defined()) || (!running_mean.defined() && !running_var.defined() && training)) && detail::getCUDAHooks().compiledWithMIOpen() - && getenv("DISABLE_MIOPEN") == NULL + && MIOPEN_ENABLED ); if (use_miopen) { From a84d5b8ef89db4d7c41d1d3789d2f7f3745a0af1 Mon Sep 17 00:00:00 2001 From: Michael Wootton Date: Tue, 25 Sep 2018 17:57:58 -0500 Subject: [PATCH 5/6] Limit MIOpen batchnorm to same-precision --- aten/src/ATen/native/Normalization.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/aten/src/ATen/native/Normalization.cpp b/aten/src/ATen/native/Normalization.cpp index 4b39dcd46e981f..18d9762dc05884 100644 --- a/aten/src/ATen/native/Normalization.cpp +++ b/aten/src/ATen/native/Normalization.cpp @@ -72,6 +72,7 @@ Tensor batch_norm( bool use_miopen = (input.type().is_cuda() && input.dim() < MIOPEN_DIM_MAX && input.type().scalarType() != at::kDouble + && (input.type().scalarType() == weight.type().scalarType()) && weight.defined() && bias.defined() && ((running_mean.defined() && running_var.defined()) || (!running_mean.defined() && !running_var.defined() && training)) From e08fdac75cd3291fdffea0a44a00eb4b7f81d778 Mon Sep 17 00:00:00 2001 From: Michael Wootton Date: Wed, 26 Sep 2018 13:57:35 -0500 Subject: [PATCH 6/6] Mark static const things actually const --- aten/src/ATen/native/Convolution.cpp | 4 ++-- aten/src/ATen/native/Normalization.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/aten/src/ATen/native/Convolution.cpp b/aten/src/ATen/native/Convolution.cpp index f8f6dd227885ce..1b8e68a1bd8ffe 100644 --- a/aten/src/ATen/native/Convolution.cpp +++ b/aten/src/ATen/native/Convolution.cpp @@ -3,8 +3,8 @@ #include "ATen/Config.h" -static int MIOPEN_DIM_MAX = 4; -static bool MIOPEN_ENABLED = getenv("DISABLE_MIOPEN") != NULL; +static const int MIOPEN_DIM_MAX = 4; +static const bool MIOPEN_ENABLED = getenv("DISABLE_MIOPEN") != NULL; namespace at { namespace native { diff --git a/aten/src/ATen/native/Normalization.cpp b/aten/src/ATen/native/Normalization.cpp index 18d9762dc05884..cb37465dab1a56 100644 --- a/aten/src/ATen/native/Normalization.cpp +++ b/aten/src/ATen/native/Normalization.cpp @@ -7,8 +7,8 @@ #include -static int MIOPEN_DIM_MAX = 4; -static bool MIOPEN_ENABLED = getenv("DISABLE_MIOPEN") != NULL; +static const int MIOPEN_DIM_MAX = 4; +static const bool MIOPEN_ENABLED = getenv("DISABLE_MIOPEN") != NULL; namespace at { namespace native {