From 11a9248d01fe2c1e946785ca9a2c5e393f549a1c Mon Sep 17 00:00:00 2001 From: rohithkrn Date: Fri, 7 Dec 2018 17:23:49 -0800 Subject: [PATCH] Enable fp16 for MIOPEN operators in Caffe2 (#14905) Summary: This PR enables fp16 MIOPEN operators in Caffe2. Pull Request resolved: https://github.com/pytorch/pytorch/pull/14905 Differential Revision: D13383439 Pulled By: bddppq fbshipit-source-id: 840afa8d08bef2952ca0039dee2423f1542bb330 --- caffe2/operators/hip/activation_ops_miopen.h | 4 ++-- caffe2/operators/hip/conv_op_miopen.hip | 19 ++++++++++++++++++- caffe2/operators/hip/elu_op_miopen.hip | 4 ++-- .../hip/local_response_normalization_op_miopen.hip | 4 ++++ caffe2/operators/hip/pool_op_miopen.hip | 4 ++++ caffe2/operators/hip/spatial_batch_norm_op_miopen.hip | 4 ++++ caffe2/python/operator_test/activation_ops_test.py | 11 +++++++---- 7 files changed, 41 insertions(+), 9 deletions(-) diff --git a/caffe2/operators/hip/activation_ops_miopen.h b/caffe2/operators/hip/activation_ops_miopen.h index af85db8..acd9db9 100644 --- a/caffe2/operators/hip/activation_ops_miopen.h +++ b/caffe2/operators/hip/activation_ops_miopen.h @@ -45,7 +45,7 @@ class MIOPENActivationOp final : public MIOPENActivationOpBase { } bool RunOnDevice() override { - return DispatchHelper>::call(this, Input(0)); + return DispatchHelper>::call(this, Input(0)); } template @@ -100,7 +100,7 @@ class MIOPENActivationGradientOp final : public MIOPENActivationOpBase { } bool RunOnDevice() override { - return DispatchHelper>::call(this, Input(0)); + return DispatchHelper>::call(this, Input(0)); } template diff --git a/caffe2/operators/hip/conv_op_miopen.hip b/caffe2/operators/hip/conv_op_miopen.hip index f42ee4a..b82a6c8 100644 --- a/caffe2/operators/hip/conv_op_miopen.hip +++ b/caffe2/operators/hip/conv_op_miopen.hip @@ -370,8 +370,15 @@ bool MIOPENConvOp::RunOnDevice() { float, // B float, // Math float>(); // Y + } else if (Input(0).IsType()) { + return DoRunWithType< + at::Half, // X + at::Half, // W + at::Half, // B + at::Half, // Math + at::Half>(); // Y } else { - LOG(FATAL) << "Only float (32bit) is supported by " + LOG(FATAL) << "Only float (32bit) and Half are supported by " << "miopen convolution, but input " << debug_def().input(0) << " has [" << Input(0).meta().name() << "]"; } @@ -621,6 +628,16 @@ bool MIOPENConvGradientOp::RunOnDevice() { float, // dX float, // dW float>(); // db + } else if (Input(0).IsType()){ + return DoRunWithType< + at::Half, // X + at::Half, // dY + at::Half, // W + at::Half, // b + at::Half, // Math + at::Half, // dX + at::Half, // dW + at::Half>(); // db } else { LOG(FATAL) << "Unsupported input types"; } diff --git a/caffe2/operators/hip/elu_op_miopen.hip b/caffe2/operators/hip/elu_op_miopen.hip index e39d2a0..fc26b8e 100644 --- a/caffe2/operators/hip/elu_op_miopen.hip +++ b/caffe2/operators/hip/elu_op_miopen.hip @@ -22,7 +22,7 @@ class MIOPENActivationOp final } bool RunOnDevice() override { - return DispatchHelper>::call(this, Input(0)); + return DispatchHelper>::call(this, Input(0)); } template @@ -85,7 +85,7 @@ class MIOPENActivationGradientOp final } bool RunOnDevice() override { - return DispatchHelper>::call(this, Input(0)); + return DispatchHelper>::call(this, Input(0)); } template diff --git a/caffe2/operators/hip/local_response_normalization_op_miopen.hip b/caffe2/operators/hip/local_response_normalization_op_miopen.hip index dfa4c32..3f6212c 100644 --- a/caffe2/operators/hip/local_response_normalization_op_miopen.hip +++ b/caffe2/operators/hip/local_response_normalization_op_miopen.hip @@ -158,6 +158,8 @@ bool MIOPEN_LRNOP::RunOnDevice() { if (X.IsType()) { return DoRunWithType(); + } else if (X.IsType()) { + return DoRunWithType(); } else { CAFFE_THROW("Unsupported input type"); } @@ -234,6 +236,8 @@ bool MIOPENLRNGradientOp::RunOnDevice() { if (dY.IsType()) { return DoRunWithType(); + } else if (dY.IsType()) { + return DoRunWithType(); } else { CAFFE_THROW("Unsupported input type"); } diff --git a/caffe2/operators/hip/pool_op_miopen.hip b/caffe2/operators/hip/pool_op_miopen.hip index 471649b..614b6cf 100644 --- a/caffe2/operators/hip/pool_op_miopen.hip +++ b/caffe2/operators/hip/pool_op_miopen.hip @@ -117,6 +117,8 @@ class MIOPENPoolOp : public ConvPoolOpBase { // TODO enable fp16 if (X.IsType()) { return DoRunWithType(); + } else if (X.IsType()){ + return DoRunWithType(); } else { LOG(FATAL) << "Unsupported input types"; } @@ -289,6 +291,8 @@ class MIOPENPoolGradientOp : public ConvPoolOpBase { if (X.IsType()) { return DoRunWithType(); + } else if (X.IsType()) { + return DoRunWithType(); } else { LOG(FATAL) << "Unsupported input types"; } diff --git a/caffe2/operators/hip/spatial_batch_norm_op_miopen.hip b/caffe2/operators/hip/spatial_batch_norm_op_miopen.hip index ecd1a4c..d833e14 100644 --- a/caffe2/operators/hip/spatial_batch_norm_op_miopen.hip +++ b/caffe2/operators/hip/spatial_batch_norm_op_miopen.hip @@ -252,6 +252,8 @@ bool MIOpenSpatialBNOp::RunOnDevice() { } if (Input(0).IsType()) { return DoRunWithType(); + } else if (Input(0).IsType()){ + return DoRunWithType(); } else { LOG(FATAL) << "Unsupported input types"; } @@ -336,6 +338,8 @@ bool MIOpenSpatialBNGradientOp::RunOnDevice() { } if (Input(0).IsType()) { return DoRunWithType(); + } else if (Input(0).IsType()) { + return DoRunWithType(); } else { LOG(FATAL) << "Unsupported input types"; } diff --git a/caffe2/python/operator_test/activation_ops_test.py b/caffe2/python/operator_test/activation_ops_test.py index 5be8b68..da1b6aa 100644 --- a/caffe2/python/operator_test/activation_ops_test.py +++ b/caffe2/python/operator_test/activation_ops_test.py @@ -5,7 +5,7 @@ from __future__ import unicode_literals import numpy as np -from hypothesis import given +from hypothesis import given, assume import hypothesis.strategies as st from caffe2.python import core, workspace @@ -41,11 +41,14 @@ class TestActivations(serial.SerializedTestCase): self.assertDeviceChecks(dc, op, [X], [0]) self.assertGradientChecks(gc, op, [X], 0, [0]) - @unittest.skipIf(not workspace.has_gpu_support, + @unittest.skipIf(not workspace.has_gpu_support and + not workspace.has_hip_support, "Relu for float16 can only run on GPU now.") @given(X=hu.tensor(dtype=np.float16), in_place=st.booleans(), - engine=st.sampled_from(["", "CUDNN"]), **hu.gcs_gpu_only) + engine=st.sampled_from(["", "CUDNN"]), **hu.gcs) def test_relu_fp16(self, X, in_place, engine, gc, dc): + # fp16 is only supported on CUDA/HIP + assume(core.IsGPUDeviceType(gc.device_type)) op = core.CreateOperator( "Relu", ["X"], @@ -68,7 +71,7 @@ class TestActivations(serial.SerializedTestCase): X[X == 0.0] += 0.02 self.assertReferenceChecks( - hu.gpu_do, + gc, op, [X], relu_ref, -- 2.7.4