From 7c441d51889d38af1949a86954dd12d74790e5da Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Thu, 13 May 2021 11:15:54 +0800 Subject: [PATCH 01/13] porting mmcv for hip --- mmcv/ops/csrc/carafe_cuda_kernel.cuh | 22 ++++++++++++++++++++-- mmcv/ops/csrc/pytorch/info.cpp | 6 ++++++ mmcv/utils/parrots_wrapper.py | 10 +++++++++- setup.py | 27 +++++++++++++++++++++++++-- 4 files changed, 60 insertions(+), 5 deletions(-) diff --git a/mmcv/ops/csrc/carafe_cuda_kernel.cuh b/mmcv/ops/csrc/carafe_cuda_kernel.cuh index e9b569d3b5..684dfbc956 100644 --- a/mmcv/ops/csrc/carafe_cuda_kernel.cuh +++ b/mmcv/ops/csrc/carafe_cuda_kernel.cuh @@ -7,7 +7,11 @@ #include "pytorch_cuda_helper.hpp" #endif +#ifdef HIP_DIFF +#define WARP_SIZE 64 +#else #define WARP_SIZE 32 +#endif #define THREADS_PER_PIXEL 32 #define MAX_SHARED_MEMORY 49152 #define MAX_SHARED_SCALAR_T 6144 // 49152 / 8 = 6144 @@ -24,6 +28,7 @@ __device__ inline int Loc2Index(const int n, const int c, const int h, int index = w + (h + (c + n * channel_num) * height) * width; return index; } +#ifndef HIP_DIFF /* TODO: move this to a common place */ template __device__ inline scalar_t min(scalar_t a, scalar_t b) { @@ -34,19 +39,28 @@ template __device__ inline scalar_t max(scalar_t a, scalar_t b) { return a > b ? a : b; } - +#endif template __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { for (int offset = 16; offset > 0; offset /= 2) - val += __shfl_down_sync(FULL_MASK, val, offset); +#ifdef HIP_DIFF + val += __shfl_down(FULL_MASK, val, offset); +#else + val += __shfl_down_sync(FULL_MASK, val, offset);#endif + return val; } template <> __device__ __forceinline__ phalf warpReduceSum(phalf val) { for (int offset = 16; offset > 0; offset /= 2) +#ifdef HIP_DIFF + __PHALF(val) += + __shfl_down(FULL_MASK, val, offset); +#else __PHALF(val) += __shfl_down_sync(FULL_MASK, static_cast<__half>(__PHALF(val)), offset); +#endif return val; } @@ -302,7 +316,11 @@ __global__ void CARAFEBackward_Mask(const int num_kernels, output_val += top_diff[top_id] * bottom_data[bottom_id]; } } +#ifdef HIP_DIFF + __syncthreads(); +#else __syncwarp(); +#endif output_val = warpReduceSum(output_val); if (lane_id == 0) { const int mask_id = diff --git a/mmcv/ops/csrc/pytorch/info.cpp b/mmcv/ops/csrc/pytorch/info.cpp index a2ebafa843..fd01c2e371 100644 --- a/mmcv/ops/csrc/pytorch/info.cpp +++ b/mmcv/ops/csrc/pytorch/info.cpp @@ -3,12 +3,15 @@ #include "pytorch_cpp_helper.hpp" #ifdef MMCV_WITH_CUDA +#ifndef HIP_DIFF #include int get_cudart_version() { return CUDART_VERSION; } #endif +#endif std::string get_compiling_cuda_version() { #ifdef MMCV_WITH_CUDA +#ifndef HIP_DIFF std::ostringstream oss; // copied from // https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/cuda/detail/CUDAHooks.cpp#L231 @@ -20,6 +23,9 @@ std::string get_compiling_cuda_version() { }; printCudaStyleVersion(get_cudart_version()); return oss.str(); +#else + return std::string("rocm not vailable"); +#endif #else return std::string("not available"); #endif diff --git a/mmcv/utils/parrots_wrapper.py b/mmcv/utils/parrots_wrapper.py index 25761be835..60c138297d 100644 --- a/mmcv/utils/parrots_wrapper.py +++ b/mmcv/utils/parrots_wrapper.py @@ -4,12 +4,20 @@ TORCH_VERSION = torch.__version__ +is_rocm_pytorch = False +if torch.__version__ >= '1.5': + from torch.utils.cpp_extension import ROCM_HOME + is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False def _get_cuda_home(): if TORCH_VERSION == 'parrots': from parrots.utils.build_extension import CUDA_HOME else: - from torch.utils.cpp_extension import CUDA_HOME + if is_rocm_pytorch: + from torch.utils.cpp_extension import ROCM_HOME + CUDA_HOME = ROCM_HOME + else: + from torch.utils.cpp_extension import CUDA_HOME return CUDA_HOME diff --git a/setup.py b/setup.py index 2b1ae6dc18..72a2b2e318 100644 --- a/setup.py +++ b/setup.py @@ -220,18 +220,41 @@ def get_extensions(): define_macros = [] extra_compile_args = {'cxx': []} - if torch.cuda.is_available() or os.getenv('FORCE_CUDA', '0') == '1': + is_rocm_pytorch = False + if torch.__version__ >= '1.5': + from torch.utils.cpp_extension import ROCM_HOME + is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False + + this_dir = "mmcv/ops/csrc/" + if is_rocm_pytorch: + from torch.utils.hipify import hipify_python + hipify_python.hipify( + project_directory=this_dir, + output_directory=this_dir, + includes="mmcv/ops/csrc/*", + show_detailed=True, + is_pytorch_extension=True, + ) + define_macros += [('MMCV_WITH_CUDA', None)] + define_macros += [('HIP_DIFF', None)] + extra_compile_args['nvcc'] = [cuda_args] if cuda_args else [] + op_files = glob.glob('./mmcv/ops/csrc/pytorch/hip/*') + extension = CUDAExtension + include_path = os.path.abspath('./mmcv/ops/csrc/hip') + + elif torch.cuda.is_available() or os.getenv('FORCE_CUDA', '0') == '1': define_macros += [('MMCV_WITH_CUDA', None)] cuda_args = os.getenv('MMCV_CUDA_ARGS') extra_compile_args['nvcc'] = [cuda_args] if cuda_args else [] op_files = glob.glob('./mmcv/ops/csrc/pytorch/*') extension = CUDAExtension + include_path = os.path.abspath('./mmcv/ops/csrc') else: print(f'Compiling {ext_name} without CUDA') op_files = glob.glob('./mmcv/ops/csrc/pytorch/*.cpp') extension = CppExtension + include_path = os.path.abspath('./mmcv/ops/csrc') - include_path = os.path.abspath('./mmcv/ops/csrc') ext_ops = extension( name=ext_name, sources=op_files, From 18a9569c97a2e48ea07a285016326d3423ee4940 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Thu, 13 May 2021 15:44:36 +0800 Subject: [PATCH 02/13] add nvcc --- mmcv/ops/csrc/carafe_cuda_kernel.cuh | 2 +- setup.py | 13 ++++++++----- 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/mmcv/ops/csrc/carafe_cuda_kernel.cuh b/mmcv/ops/csrc/carafe_cuda_kernel.cuh index 684dfbc956..84656564c1 100644 --- a/mmcv/ops/csrc/carafe_cuda_kernel.cuh +++ b/mmcv/ops/csrc/carafe_cuda_kernel.cuh @@ -47,7 +47,7 @@ __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { val += __shfl_down(FULL_MASK, val, offset); #else val += __shfl_down_sync(FULL_MASK, val, offset);#endif - +#endif return val; } diff --git a/setup.py b/setup.py index 72a2b2e318..abeb7da2d3 100644 --- a/setup.py +++ b/setup.py @@ -223,20 +223,23 @@ def get_extensions(): is_rocm_pytorch = False if torch.__version__ >= '1.5': from torch.utils.cpp_extension import ROCM_HOME - is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False + is_rocm_pytorch = True if ((torch.version.hip is not None) and + (ROCM_HOME is not None)) else False - this_dir = "mmcv/ops/csrc/" + this_dir = 'mmcv/ops/csrc/' if is_rocm_pytorch: from torch.utils.hipify import hipify_python + hipify_python.hipify( project_directory=this_dir, - output_directory=this_dir, - includes="mmcv/ops/csrc/*", - show_detailed=True, + output_directory=this_dir, + includes='mmcv/ops/csrc/*', + show_detailed=True, is_pytorch_extension=True, ) define_macros += [('MMCV_WITH_CUDA', None)] define_macros += [('HIP_DIFF', None)] + cuda_args = os.getenv('MMCV_CUDA_ARGS') extra_compile_args['nvcc'] = [cuda_args] if cuda_args else [] op_files = glob.glob('./mmcv/ops/csrc/pytorch/hip/*') extension = CUDAExtension From 59480dddb1f1d77451a809daaefe99c14f077038 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Thu, 13 May 2021 16:28:31 +0800 Subject: [PATCH 03/13] fix format --- mmcv/ops/csrc/carafe_cuda_kernel.cuh | 2 +- mmcv/utils/parrots_wrapper.py | 4 +++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/mmcv/ops/csrc/carafe_cuda_kernel.cuh b/mmcv/ops/csrc/carafe_cuda_kernel.cuh index 84656564c1..ec408ea6b3 100644 --- a/mmcv/ops/csrc/carafe_cuda_kernel.cuh +++ b/mmcv/ops/csrc/carafe_cuda_kernel.cuh @@ -46,7 +46,7 @@ __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { #ifdef HIP_DIFF val += __shfl_down(FULL_MASK, val, offset); #else - val += __shfl_down_sync(FULL_MASK, val, offset);#endif + val += __shfl_down_sync(FULL_MASK, val, offset); #endif return val; } diff --git a/mmcv/utils/parrots_wrapper.py b/mmcv/utils/parrots_wrapper.py index 60c138297d..a3f4eb9f84 100644 --- a/mmcv/utils/parrots_wrapper.py +++ b/mmcv/utils/parrots_wrapper.py @@ -7,7 +7,9 @@ is_rocm_pytorch = False if torch.__version__ >= '1.5': from torch.utils.cpp_extension import ROCM_HOME - is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False + is_rocm_pytorch = True if ((torch.version.hip is not None) and + (ROCM_HOME is not None)) else False + def _get_cuda_home(): if TORCH_VERSION == 'parrots': From 01d72d00ad456f5af79cada384db77b1c8d77cf6 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Fri, 14 May 2021 16:32:47 +0800 Subject: [PATCH 04/13] fix format --- mmcv/ops/csrc/carafe_cuda_kernel.cuh | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/mmcv/ops/csrc/carafe_cuda_kernel.cuh b/mmcv/ops/csrc/carafe_cuda_kernel.cuh index ec408ea6b3..fb2c770787 100644 --- a/mmcv/ops/csrc/carafe_cuda_kernel.cuh +++ b/mmcv/ops/csrc/carafe_cuda_kernel.cuh @@ -55,8 +55,7 @@ template <> __device__ __forceinline__ phalf warpReduceSum(phalf val) { for (int offset = 16; offset > 0; offset /= 2) #ifdef HIP_DIFF - __PHALF(val) += - __shfl_down(FULL_MASK, val, offset); + __PHALF(val) += __shfl_down(FULL_MASK, val, offset); #else __PHALF(val) += __shfl_down_sync(FULL_MASK, static_cast<__half>(__PHALF(val)), offset); From 80143d6461f4d4af6fc2bbfeb8d57cf27179b45a Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Wed, 7 Jul 2021 14:16:24 +0800 Subject: [PATCH 05/13] fix bug for carafe --- mmcv/ops/csrc/carafe_cuda_kernel.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/mmcv/ops/csrc/carafe_cuda_kernel.cuh b/mmcv/ops/csrc/carafe_cuda_kernel.cuh index fb2c770787..f2fe702df2 100644 --- a/mmcv/ops/csrc/carafe_cuda_kernel.cuh +++ b/mmcv/ops/csrc/carafe_cuda_kernel.cuh @@ -42,9 +42,9 @@ __device__ inline scalar_t max(scalar_t a, scalar_t b) { #endif template __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { - for (int offset = 16; offset > 0; offset /= 2) + for (int offset = WARP_SIZE/2; offset > 0; offset /= 2) #ifdef HIP_DIFF - val += __shfl_down(FULL_MASK, val, offset); + val += __shfl_down(val, offset); #else val += __shfl_down_sync(FULL_MASK, val, offset); #endif @@ -53,7 +53,7 @@ __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { template <> __device__ __forceinline__ phalf warpReduceSum(phalf val) { - for (int offset = 16; offset > 0; offset /= 2) + for (int offset = WARP_SIZE/2; offset > 0; offset /= 2) #ifdef HIP_DIFF __PHALF(val) += __shfl_down(FULL_MASK, val, offset); #else From 40cc3f58c48001b9c99473b50bcae776de852a9c Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Wed, 7 Jul 2021 16:43:59 +0800 Subject: [PATCH 06/13] fix test_utils because rocm_torch not allow set torch.backends.cudnn.benchmark to false --- tests/test_runner/test_utils.py | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/tests/test_runner/test_utils.py b/tests/test_runner/test_utils.py index 3983e80cd7..56e586a777 100644 --- a/tests/test_runner/test_utils.py +++ b/tests/test_runner/test_utils.py @@ -6,6 +6,12 @@ from mmcv.runner import set_random_seed +is_rocm_pytorch = False +if torch.__version__ >= '1.5': + from torch.utils.cpp_extension import ROCM_HOME + is_rocm_pytorch = True if ((torch.version.hip is not None) and + (ROCM_HOME is not None)) else False + def test_set_random_seed(): set_random_seed(0) @@ -21,7 +27,10 @@ def test_set_random_seed(): b_np_random = np.random.rand(2, 2) b_torch_random = torch.rand(2, 2) assert torch.backends.cudnn.deterministic is True - assert torch.backends.cudnn.benchmark is False + if is_rocm_pytorch: + assert torch.backends.cudnn.benchmark is True + else: + assert torch.backends.cudnn.benchmark is False assert a_random == b_random assert np.equal(a_np_random, b_np_random).all() From 7bf604ac7c29c8af71488ee7181b85a4366b98fa Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Wed, 7 Jul 2021 17:48:48 +0800 Subject: [PATCH 07/13] add LOOSEVERSION --- mmcv/utils/parrots_wrapper.py | 3 ++- setup.py | 5 ++++- tests/test_runner/test_utils.py | 4 +++- 3 files changed, 9 insertions(+), 3 deletions(-) diff --git a/mmcv/utils/parrots_wrapper.py b/mmcv/utils/parrots_wrapper.py index fc8ed516c7..ed56b3958e 100644 --- a/mmcv/utils/parrots_wrapper.py +++ b/mmcv/utils/parrots_wrapper.py @@ -1,3 +1,4 @@ +from distutils.version import LooseVersion from functools import partial import torch @@ -5,7 +6,7 @@ TORCH_VERSION = torch.__version__ is_rocm_pytorch = False -if torch.__version__ >= '1.5': +if (LooseVersion(TORCH_VERSION) >= LooseVersion('1.5')): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/setup.py b/setup.py index abeb7da2d3..dfe704c297 100644 --- a/setup.py +++ b/setup.py @@ -1,9 +1,12 @@ import glob import os import re +from distutils.version import LooseVersion from pkg_resources import DistributionNotFound, get_distribution from setuptools import find_packages, setup +from mmcv.utils import TORCH_VERSION + EXT_TYPE = '' try: import torch @@ -221,7 +224,7 @@ def get_extensions(): extra_compile_args = {'cxx': []} is_rocm_pytorch = False - if torch.__version__ >= '1.5': + if (LooseVersion(TORCH_VERSION) >= LooseVersion('1.5')): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/tests/test_runner/test_utils.py b/tests/test_runner/test_utils.py index 56e586a777..05e74f368d 100644 --- a/tests/test_runner/test_utils.py +++ b/tests/test_runner/test_utils.py @@ -1,13 +1,15 @@ import os import random +from distutils.version import LooseVersion import numpy as np import torch from mmcv.runner import set_random_seed +from mmcv.utils import TORCH_VERSION is_rocm_pytorch = False -if torch.__version__ >= '1.5': +if (LooseVersion(TORCH_VERSION) >= LooseVersion('1.5')): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False From d6ffb3637cf0255fe1c950b038d2bf3d04b9358e Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Wed, 7 Jul 2021 17:53:50 +0800 Subject: [PATCH 08/13] fix format --- mmcv/utils/parrots_wrapper.py | 2 +- setup.py | 2 +- tests/test_runner/test_utils.py | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/mmcv/utils/parrots_wrapper.py b/mmcv/utils/parrots_wrapper.py index ed56b3958e..2a06edde77 100644 --- a/mmcv/utils/parrots_wrapper.py +++ b/mmcv/utils/parrots_wrapper.py @@ -6,7 +6,7 @@ TORCH_VERSION = torch.__version__ is_rocm_pytorch = False -if (LooseVersion(TORCH_VERSION) >= LooseVersion('1.5')): +if LooseVersion(TORCH_VERSION) >= LooseVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/setup.py b/setup.py index dfe704c297..acd5520f62 100644 --- a/setup.py +++ b/setup.py @@ -224,7 +224,7 @@ def get_extensions(): extra_compile_args = {'cxx': []} is_rocm_pytorch = False - if (LooseVersion(TORCH_VERSION) >= LooseVersion('1.5')): + if LooseVersion(TORCH_VERSION) >= LooseVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/tests/test_runner/test_utils.py b/tests/test_runner/test_utils.py index 05e74f368d..851455daf7 100644 --- a/tests/test_runner/test_utils.py +++ b/tests/test_runner/test_utils.py @@ -9,7 +9,7 @@ from mmcv.utils import TORCH_VERSION is_rocm_pytorch = False -if (LooseVersion(TORCH_VERSION) >= LooseVersion('1.5')): +if LooseVersion(TORCH_VERSION) >= LooseVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False From d85af3a07df0d465e2ebe7b144e2d6841b0a8f29 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Thu, 8 Jul 2021 14:17:11 +0800 Subject: [PATCH 09/13] fix format of version --- mmcv/utils/parrots_wrapper.py | 4 ++-- setup.py | 4 ++-- tests/test_runner/test_utils.py | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/mmcv/utils/parrots_wrapper.py b/mmcv/utils/parrots_wrapper.py index 2a06edde77..99e1b8aa41 100644 --- a/mmcv/utils/parrots_wrapper.py +++ b/mmcv/utils/parrots_wrapper.py @@ -1,4 +1,4 @@ -from distutils.version import LooseVersion +from distutils.version import StrictVersion from functools import partial import torch @@ -6,7 +6,7 @@ TORCH_VERSION = torch.__version__ is_rocm_pytorch = False -if LooseVersion(TORCH_VERSION) >= LooseVersion('1.5'): +if StrictVersion(TORCH_VERSION) >= StrictVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/setup.py b/setup.py index acd5520f62..470ba7c0b5 100644 --- a/setup.py +++ b/setup.py @@ -1,7 +1,7 @@ import glob import os import re -from distutils.version import LooseVersion +from distutils.version import StrictVersion from pkg_resources import DistributionNotFound, get_distribution from setuptools import find_packages, setup @@ -224,7 +224,7 @@ def get_extensions(): extra_compile_args = {'cxx': []} is_rocm_pytorch = False - if LooseVersion(TORCH_VERSION) >= LooseVersion('1.5'): + if StrictVersion(TORCH_VERSION) >= StrictVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/tests/test_runner/test_utils.py b/tests/test_runner/test_utils.py index 851455daf7..596f7aeb81 100644 --- a/tests/test_runner/test_utils.py +++ b/tests/test_runner/test_utils.py @@ -1,6 +1,6 @@ import os import random -from distutils.version import LooseVersion +from distutils.version import StrictVersion import numpy as np import torch @@ -9,7 +9,7 @@ from mmcv.utils import TORCH_VERSION is_rocm_pytorch = False -if LooseVersion(TORCH_VERSION) >= LooseVersion('1.5'): +if StrictVersion(TORCH_VERSION) >= StrictVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False From 4b3189aa46c4aef6875484c48fedfc35439e8625 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Thu, 8 Jul 2021 17:38:17 +0800 Subject: [PATCH 10/13] fix code format --- mmcv/ops/csrc/carafe_cuda_kernel.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mmcv/ops/csrc/carafe_cuda_kernel.cuh b/mmcv/ops/csrc/carafe_cuda_kernel.cuh index f2fe702df2..4bf11694f3 100644 --- a/mmcv/ops/csrc/carafe_cuda_kernel.cuh +++ b/mmcv/ops/csrc/carafe_cuda_kernel.cuh @@ -42,7 +42,7 @@ __device__ inline scalar_t max(scalar_t a, scalar_t b) { #endif template __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { - for (int offset = WARP_SIZE/2; offset > 0; offset /= 2) + for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2) #ifdef HIP_DIFF val += __shfl_down(val, offset); #else @@ -53,7 +53,7 @@ __device__ __forceinline__ scalar_t warpReduceSum(scalar_t val) { template <> __device__ __forceinline__ phalf warpReduceSum(phalf val) { - for (int offset = WARP_SIZE/2; offset > 0; offset /= 2) + for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2) #ifdef HIP_DIFF __PHALF(val) += __shfl_down(FULL_MASK, val, offset); #else From 22f0c8e380291b0fda2f16aadf55a96acb49a8bf Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Thu, 8 Jul 2021 20:31:42 +0800 Subject: [PATCH 11/13] test for yaml --- .github/workflows/build.yml | 2 ++ 1 file changed, 2 insertions(+) diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 0b90cbf215..f5e75f3c5d 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -116,6 +116,8 @@ jobs: - name: Install Pillow run: pip install Pillow==6.2.2 if: ${{matrix.torchvision == '0.4.2'}} + - name: Install numpy and yaml + run: pip install -r requirements/runtime.txt - name: Install PyTorch run: pip install torch==${{matrix.torch}}+cpu torchvision==${{matrix.torchvision}}+cpu -f https://download.pytorch.org/whl/torch_stable.html - name: Build and install From 8c18ffce02c04e01a570201a553e610c4efa4261 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Fri, 9 Jul 2021 09:18:52 +0800 Subject: [PATCH 12/13] fix bug for citest --- .github/workflows/build.yml | 2 -- setup.py | 5 ++--- 2 files changed, 2 insertions(+), 5 deletions(-) diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index f5e75f3c5d..0b90cbf215 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -116,8 +116,6 @@ jobs: - name: Install Pillow run: pip install Pillow==6.2.2 if: ${{matrix.torchvision == '0.4.2'}} - - name: Install numpy and yaml - run: pip install -r requirements/runtime.txt - name: Install PyTorch run: pip install torch==${{matrix.torch}}+cpu torchvision==${{matrix.torchvision}}+cpu -f https://download.pytorch.org/whl/torch_stable.html - name: Build and install diff --git a/setup.py b/setup.py index 470ba7c0b5..a770686938 100644 --- a/setup.py +++ b/setup.py @@ -5,11 +5,10 @@ from pkg_resources import DistributionNotFound, get_distribution from setuptools import find_packages, setup -from mmcv.utils import TORCH_VERSION +import torch EXT_TYPE = '' try: - import torch if torch.__version__ == 'parrots': from parrots.utils.build_extension import BuildExtension EXT_TYPE = 'parrots' @@ -224,7 +223,7 @@ def get_extensions(): extra_compile_args = {'cxx': []} is_rocm_pytorch = False - if StrictVersion(TORCH_VERSION) >= StrictVersion('1.5'): + if StrictVersion(torch.__version__) >= StrictVersion('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False From e82c5e0625f28dd98cf4620ff0c68d19f4312ca7 Mon Sep 17 00:00:00 2001 From: XUANBABY Date: Fri, 9 Jul 2021 10:03:58 +0800 Subject: [PATCH 13/13] fix bug for how to get torch._version_ at setup.py --- mmcv/utils/parrots_wrapper.py | 4 ++-- setup.py | 8 +++----- tests/test_runner/test_utils.py | 4 ++-- 3 files changed, 7 insertions(+), 9 deletions(-) diff --git a/mmcv/utils/parrots_wrapper.py b/mmcv/utils/parrots_wrapper.py index 99e1b8aa41..4cb4be98f2 100644 --- a/mmcv/utils/parrots_wrapper.py +++ b/mmcv/utils/parrots_wrapper.py @@ -1,12 +1,12 @@ -from distutils.version import StrictVersion from functools import partial +from pkg_resources import parse_version import torch TORCH_VERSION = torch.__version__ is_rocm_pytorch = False -if StrictVersion(TORCH_VERSION) >= StrictVersion('1.5'): +if parse_version(TORCH_VERSION) >= parse_version('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/setup.py b/setup.py index a770686938..4122c814d2 100644 --- a/setup.py +++ b/setup.py @@ -1,14 +1,12 @@ import glob import os import re -from distutils.version import StrictVersion -from pkg_resources import DistributionNotFound, get_distribution +from pkg_resources import DistributionNotFound, get_distribution, parse_version from setuptools import find_packages, setup -import torch - EXT_TYPE = '' try: + import torch if torch.__version__ == 'parrots': from parrots.utils.build_extension import BuildExtension EXT_TYPE = 'parrots' @@ -223,7 +221,7 @@ def get_extensions(): extra_compile_args = {'cxx': []} is_rocm_pytorch = False - if StrictVersion(torch.__version__) >= StrictVersion('1.5'): + if parse_version(torch.__version__) >= parse_version('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False diff --git a/tests/test_runner/test_utils.py b/tests/test_runner/test_utils.py index 596f7aeb81..88e0629c28 100644 --- a/tests/test_runner/test_utils.py +++ b/tests/test_runner/test_utils.py @@ -1,6 +1,6 @@ import os import random -from distutils.version import StrictVersion +from pkg_resources import parse_version import numpy as np import torch @@ -9,7 +9,7 @@ from mmcv.utils import TORCH_VERSION is_rocm_pytorch = False -if StrictVersion(TORCH_VERSION) >= StrictVersion('1.5'): +if parse_version(TORCH_VERSION) >= parse_version('1.5'): from torch.utils.cpp_extension import ROCM_HOME is_rocm_pytorch = True if ((torch.version.hip is not None) and (ROCM_HOME is not None)) else False