[arch-commits] Commit in python-pytorch/trunk (9 files)
Konstantin Gizdov
kgizdov at gemini.archlinux.org
Sun Oct 31 22:49:00 UTC 2021
Date: Sunday, October 31, 2021 @ 22:49:00
Author: kgizdov
Revision: 1035135
upgpkg: python-pytorch 1.10.0-1: upgpkg: python-pytorch 1.10.0-1
Added:
python-pytorch/trunk/66219.patch
python-pytorch/trunk/fix-jit-frontend-nullptr-deref.patch
python-pytorch/trunk/fix_old_nnapi_lite_interpreter_config.patch
Modified:
python-pytorch/trunk/PKGBUILD
python-pytorch/trunk/fix-building-for-torchvision.patch
python-pytorch/trunk/fix_c10.patch
Deleted:
python-pytorch/trunk/benchmark-gcc11.patch
python-pytorch/trunk/disable_non_x86_64.patch
python-pytorch/trunk/xnnpack-gcc11.patch
---------------------------------------------+
66219.patch | 3002 ++++++++++++++++++++++++++
PKGBUILD | 135 -
benchmark-gcc11.patch | 30
disable_non_x86_64.patch | 15
fix-building-for-torchvision.patch | 6
fix-jit-frontend-nullptr-deref.patch | 12
fix_c10.patch | 2
fix_old_nnapi_lite_interpreter_config.patch | 33
xnnpack-gcc11.patch | 33
9 files changed, 3115 insertions(+), 153 deletions(-)
Added: 66219.patch
===================================================================
--- 66219.patch (rev 0)
+++ 66219.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -0,0 +1,3002 @@
+From 55473c81535db8890d35e29cff852b737954ce80 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Wed, 6 Oct 2021 11:52:58 -0700
+Subject: [PATCH 01/30] Refactor cub namespace
+
+---
+ aten/src/ATen/cuda/cub.cuh | 102 +++++++++---------
+ aten/src/ATen/cuda/cub_definitions.cuh | 17 +++
+ caffe2/core/context_gpu.cu | 1 +
+ caffe2/operators/accuracy_op.cu | 1 +
+ caffe2/operators/affine_channel_op.cu | 1 +
+ caffe2/operators/arg_ops.cu | 2 +-
+ caffe2/operators/batch_moments_op.cu | 1 +
+ caffe2/operators/batch_sparse_to_dense_op.cu | 1 +
+ caffe2/operators/boolean_mask_ops.cu | 2 +-
+ caffe2/operators/cross_entropy_op.cu | 1 +
+ caffe2/operators/distance_op.cu | 1 +
+ caffe2/operators/elementwise_div_op.cu | 2 +-
+ caffe2/operators/elementwise_linear_op.cu | 1 +
+ caffe2/operators/elementwise_mul_op.cu | 2 +-
+ caffe2/operators/elementwise_ops.cu | 1 +
+ caffe2/operators/find_op.cu | 1 +
+ caffe2/operators/generate_proposals_op.cu | 2 +-
+ caffe2/operators/normalize_ops.cu | 1 +
+ caffe2/operators/one_hot_ops.cu | 1 +
+ caffe2/operators/pack_segments.cu | 2 +-
+ caffe2/operators/prelu_op.cu | 1 +
+ caffe2/operators/reduce_front_back_max_ops.cu | 1 +
+ .../reduce_front_back_sum_mean_ops.cu | 1 +
+ caffe2/operators/reduction_ops.cu | 2 +-
+ caffe2/operators/rmac_regions_op.cu | 7 ++
+ caffe2/operators/segment_reduction_op_gpu.cuh | 1 +
+ caffe2/operators/sequence_ops.cu | 2 +-
+ caffe2/operators/softmax_ops.cu | 1 +
+ .../operators/spatial_batch_norm_op_impl.cuh | 2 +-
+ caffe2/sgd/adagrad_fused_op_gpu.cu | 1 +
+ caffe2/sgd/adagrad_op_gpu.cu | 1 +
+ caffe2/sgd/adam_op_gpu.cu | 1 +
+ caffe2/utils/math/reduce.cu | 2 +-
+ caffe2/utils/math/reduce.cuh | 2 +-
+ caffe2/utils/math_gpu.cu | 2 +-
+ cmake/Dependencies.cmake | 5 +
+ 36 files changed, 115 insertions(+), 60 deletions(-)
+ create mode 100644 aten/src/ATen/cuda/cub_definitions.cuh
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 5d8ae777ebef..39938efc48be 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -5,16 +5,24 @@
+ #include <iterator>
+ #include <limits>
+
+-// include cub in a safe manner, see:
+-// https://github.com/pytorch/pytorch/pull/55292
++#include <ATen/cuda/cub_definitions.cuh>
++
++#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++
++#include <cub/cub.cuh>
++
++#else
++
+ #undef CUB_NS_POSTFIX //undef to avoid redefinition warnings
+ #undef CUB_NS_PREFIX
+-#define CUB_NS_PREFIX namespace at { namespace cuda { namespace detail {
+-#define CUB_NS_POSTFIX }}}
++#define CUB_NS_PREFIX namespace at_cuda_detail {
++#define CUB_NS_POSTFIX }
+ #include <cub/cub.cuh>
+ #undef CUB_NS_POSTFIX
+ #undef CUB_NS_PREFIX
+
++#endif
++
+ #include <ATen/cuda/Exceptions.h>
+ #include <c10/cuda/CUDACachingAllocator.h>
+ #include <c10/cuda/CUDAStream.h>
+@@ -33,16 +41,41 @@
+ #define NO_ROCM(x)
+ #else
+ #define NO_ROCM(x) x
++#endif
+
+-namespace at { namespace native {
++#if !defined(USE_ROCM) && !CUB_SUPPORTS_NV_BFLOAT16()
+
+-namespace cub = at::cuda::detail::cub;
++namespace at_cuda_detail {
++// backport https://github.com/NVIDIA/cub/pull/306 for c10::BFloat16
+
+-}}
++template <>
++struct cub::FpLimits<c10::BFloat16>
++{
++ static __host__ __device__ __forceinline__ c10::BFloat16 Max() {
++ unsigned short max_word = 0x7F7F;
++ return reinterpret_cast<c10::BFloat16&>(max_word);
++ }
++
++ static __host__ __device__ __forceinline__ c10::BFloat16 Lowest() {
++ unsigned short lowest_word = 0xFF7F;
++ return reinterpret_cast<c10::BFloat16&>(lowest_word);
++ }
++};
++
++template <> struct cub::NumericTraits<c10::BFloat16>: cub::BaseTraits<cub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
++}
+ #endif
+
++namespace at { namespace native {
++namespace cub = at_cuda_detail::cub;
++}}
++namespace caffe2 {
++namespace cub = at_cuda_detail::cub;
++}
++
+ namespace at {
+ namespace cuda {
++namespace cub {
+
+ namespace detail {
+
+@@ -55,44 +88,17 @@ struct cuda_type<c10::Half> {
+ using type = __half;
+ };
+
+-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11050
+-// cub sort support for __nv_bfloat16 is added to cub 1.13 in
+-// https://github.com/NVIDIA/cub/pull/306 and according to
+-// https://github.com/NVIDIA/cub#releases, 1.13 is included in
+-// CUDA Toolkit 11.5
++#if CUB_SUPPORTS_NV_BFLOAT16()
+
+-// waiting for https://github.com/NVIDIA/cub/pull/306 to land on CUDA
+ template<>
+ struct cuda_type<c10::BFloat16> {
+ using type = __nv_bfloat16;
+ };
+
+-#elif !defined(__HIP_PLATFORM_HCC__)
+-
+-// backport https://github.com/NVIDIA/cub/pull/306 for c10::BFloat16
+-
+-template <>
+-struct cub::FpLimits<c10::BFloat16>
+-{
+- static __host__ __device__ __forceinline__ c10::BFloat16 Max() {
+- unsigned short max_word = 0x7F7F;
+- return reinterpret_cast<c10::BFloat16&>(max_word);
+- }
+-
+- static __host__ __device__ __forceinline__ c10::BFloat16 Lowest() {
+- unsigned short lowest_word = 0xFF7F;
+- return reinterpret_cast<c10::BFloat16&>(lowest_word);
+- }
+-};
+-
+-template <> struct cub::NumericTraits<c10::BFloat16>: cub::BaseTraits<cub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
+-
+ #endif
+
+ } // namespace detail
+
+-namespace cub {
+-
+ inline int get_num_bits(uint64_t max_key) {
+ int num_bits = 1;
+ while (max_key > 1) {
+@@ -115,11 +121,11 @@ static inline void sort_keys(
+ key_t_ *keys_out_ = reinterpret_cast<key_t_*>(keys_out);
+
+ if (descending) {
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceRadixSort::SortKeysDescending,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceRadixSort::SortKeysDescending,
+ keys_in_, keys_out_, n,
+ begin_bit, end_bit, c10::cuda::getCurrentCUDAStream());
+ } else {
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceRadixSort::SortKeys,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceRadixSort::SortKeys,
+ keys_in_, keys_out_, n,
+ begin_bit, end_bit, c10::cuda::getCurrentCUDAStream());
+ }
+@@ -147,11 +153,11 @@ static inline void sort_pairs(
+ key_t_ *keys_out_ = reinterpret_cast<key_t_*>(keys_out);
+
+ if (descending) {
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceRadixSort::SortPairsDescending,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceRadixSort::SortPairsDescending,
+ keys_in_, keys_out_, values_in, values_out, n,
+ begin_bit, end_bit, c10::cuda::getCurrentCUDAStream());
+ } else {
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceRadixSort::SortPairs,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceRadixSort::SortPairs,
+ keys_in_, keys_out_, values_in, values_out, n,
+ begin_bit, end_bit, c10::cuda::getCurrentCUDAStream());
+ }
+@@ -183,12 +189,12 @@ static inline void segmented_sort_pairs(
+ key_t_ *keys_out_ = reinterpret_cast<key_t_*>(keys_out);
+
+ if (descending) {
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceSegmentedRadixSort::SortPairsDescending,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceSegmentedRadixSort::SortPairsDescending,
+ keys_in_, keys_out_, values_in, values_out,
+ num_elements, num_segments, begin_offsets, end_offsets,
+ begin_bit, end_bit, c10::cuda::getCurrentCUDAStream());
+ } else {
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceSegmentedRadixSort::SortPairs,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceSegmentedRadixSort::SortPairs,
+ keys_in_, keys_out_, values_in, values_out,
+ num_elements, num_segments, begin_offsets, end_offsets,
+ begin_bit, end_bit, c10::cuda::getCurrentCUDAStream());
+@@ -240,7 +246,7 @@ inline void inclusive_scan(InputIteratorT input, OutputIteratorT output, ScanOpT
+ // so split at int_max/2
+ constexpr int max_cub_size = std::numeric_limits<int>::max() / 2 + 1; // 2**30
+ int size_cub = std::min<int64_t>(num_items, max_cub_size);
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceScan::InclusiveScan,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceScan::InclusiveScan,
+ input,
+ output,
+ scan_op,
+@@ -260,7 +266,7 @@ inline void inclusive_scan(InputIteratorT input, OutputIteratorT output, ScanOpT
+ first_elem_ptr,
+ scan_op);
+ C10_CUDA_KERNEL_LAUNCH_CHECK();
+- using ArgIndexInputIterator = NO_ROCM(detail)::cub::ArgIndexInputIterator<InputIteratorT>;
++ using ArgIndexInputIterator = NO_ROCM(at_cuda_detail)::cub::ArgIndexInputIterator<InputIteratorT>;
+ using tuple = typename ArgIndexInputIterator::value_type;
+ auto input_iter_transform = [=] __device__ (const tuple &x)->input_t {
+ if (x.key == 0) {
+@@ -269,9 +275,9 @@ inline void inclusive_scan(InputIteratorT input, OutputIteratorT output, ScanOpT
+ return x.value;
+ }
+ };
+- auto input_ = NO_ROCM(detail)::cub::TransformInputIterator<input_t, decltype(input_iter_transform), ArgIndexInputIterator>(
++ auto input_ = NO_ROCM(at_cuda_detail)::cub::TransformInputIterator<input_t, decltype(input_iter_transform), ArgIndexInputIterator>(
+ ArgIndexInputIterator(input + i), input_iter_transform);
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceScan::InclusiveScan,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceScan::InclusiveScan,
+ input_,
+ output + i,
+ scan_op,
+@@ -287,7 +293,7 @@ inline void exclusive_scan(InputIteratorT input, OutputIteratorT output, ScanOpT
+ // so split at int_max/2
+ constexpr int max_cub_size = std::numeric_limits<int>::max() / 2 + 1; // 2**30
+ int size_cub = std::min<int64_t>(num_items, max_cub_size);
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceScan::ExclusiveScan,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceScan::ExclusiveScan,
+ input,
+ output,
+ scan_op,
+@@ -309,7 +315,7 @@ inline void exclusive_scan(InputIteratorT input, OutputIteratorT output, ScanOpT
+ C10_CUDA_KERNEL_LAUNCH_CHECK();
+ auto input_ = impl::chained_iterator<InitValueT, InputIteratorT>{
+ input + i, first_elem_ptr};
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceScan::InclusiveScan,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceScan::InclusiveScan,
+ input_,
+ output + i,
+ scan_op,
+@@ -322,7 +328,7 @@ template<typename InputIteratorT , typename OutputIteratorT , typename NumSelect
+ inline void unique(InputIteratorT input, OutputIteratorT output, NumSelectedIteratorT num_selected_out, int64_t num_items) {
+ TORCH_CHECK(num_items <= std::numeric_limits<int>::max(),
+ "cub unique does not support more than INT_MAX elements");
+- CUB_WRAPPER(NO_ROCM(detail)::cub::DeviceSelect::Unique,
++ CUB_WRAPPER(NO_ROCM(at_cuda_detail)::cub::DeviceSelect::Unique,
+ input, output, num_selected_out, num_items, at::cuda::getCurrentCUDAStream());
+ }
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+new file mode 100644
+index 000000000000..a7694fda4168
+--- /dev/null
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -0,0 +1,17 @@
++#include <cub/version.cuh>
++
++// cub sort support for __nv_bfloat16 is added to cub 1.13 in:
++// https://github.com/NVIDIA/cub/pull/306
++#if CUB_VERSION >= 101300
++#define CUB_SUPPORTS_NV_BFLOAT16() true
++#elif
++#define CUB_SUPPORTS_NV_BFLOAT16() false
++#endif
++
++// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.14 in:
++// https://github.com/NVIDIA/cub/pull/326
++#if CUB_VERSION >= 101400
++#define CUB_SUPPORTS_WRAPPED_NAMESPACE() true
++#elif
++#define CUB_SUPPORTS_WRAPPED_NAMESPACE() false
++#endif
+\ No newline at end of file
+diff --git a/caffe2/core/context_gpu.cu b/caffe2/core/context_gpu.cu
+index c2b89945ada9..475ed61ab4f7 100644
+--- a/caffe2/core/context_gpu.cu
++++ b/caffe2/core/context_gpu.cu
+@@ -4,6 +4,7 @@
+ #include <string>
+ #include <unordered_map>
+
++#include <ATen/cuda/cub.cuh>
+ #include <ATen/Context.h>
+ #include <c10/cuda/CUDAFunctions.h>
+ #include <c10/cuda/CUDACachingAllocator.h>
+diff --git a/caffe2/operators/accuracy_op.cu b/caffe2/operators/accuracy_op.cu
+index f06663d71a90..7ad2b09c238a 100644
+--- a/caffe2/operators/accuracy_op.cu
++++ b/caffe2/operators/accuracy_op.cu
+@@ -3,6 +3,7 @@
+ #include "caffe2/utils/GpuAtomics.cuh"
+ #include "caffe2/utils/math.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/affine_channel_op.cu b/caffe2/operators/affine_channel_op.cu
+index adf4ac55c0fc..f3d9e22c7e8c 100644
+--- a/caffe2/operators/affine_channel_op.cu
++++ b/caffe2/operators/affine_channel_op.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/affine_channel_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/arg_ops.cu b/caffe2/operators/arg_ops.cu
+index 7e90d25b836b..fbefe0774376 100644
+--- a/caffe2/operators/arg_ops.cu
++++ b/caffe2/operators/arg_ops.cu
+@@ -2,8 +2,8 @@
+
+ #include <limits>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_moments_op.cu b/caffe2/operators/batch_moments_op.cu
+index 4b693b5c04e2..65c43200e5bd 100644
+--- a/caffe2/operators/batch_moments_op.cu
++++ b/caffe2/operators/batch_moments_op.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/batch_moments_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_sparse_to_dense_op.cu b/caffe2/operators/batch_sparse_to_dense_op.cu
+index aea2035a5d3d..2cb09deb8668 100644
+--- a/caffe2/operators/batch_sparse_to_dense_op.cu
++++ b/caffe2/operators/batch_sparse_to_dense_op.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/batch_sparse_to_dense_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/device/device_scan.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/boolean_mask_ops.cu b/caffe2/operators/boolean_mask_ops.cu
+index 214b7c13ba3c..c87688f51d64 100644
+--- a/caffe2/operators/boolean_mask_ops.cu
++++ b/caffe2/operators/boolean_mask_ops.cu
+@@ -3,7 +3,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/boolean_mask_ops.h"
+
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/cross_entropy_op.cu b/caffe2/operators/cross_entropy_op.cu
+index 380e80399fc3..95f3ffddbf1f 100644
+--- a/caffe2/operators/cross_entropy_op.cu
++++ b/caffe2/operators/cross_entropy_op.cu
+@@ -1,4 +1,5 @@
+ #include <assert.h>
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/distance_op.cu b/caffe2/operators/distance_op.cu
+index 3a8bb337d541..d94691d5a9d9 100644
+--- a/caffe2/operators/distance_op.cu
++++ b/caffe2/operators/distance_op.cu
+@@ -4,6 +4,7 @@
+ #include "caffe2/operators/distance_op.h"
+ #include "caffe2/utils/conversions.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_div_op.cu b/caffe2/operators/elementwise_div_op.cu
+index 42b103a0f110..ca9682326324 100644
+--- a/caffe2/operators/elementwise_div_op.cu
++++ b/caffe2/operators/elementwise_div_op.cu
+@@ -3,8 +3,8 @@
+ #include <algorithm>
+ #include <functional>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/elementwise_ops_utils.h"
+diff --git a/caffe2/operators/elementwise_linear_op.cu b/caffe2/operators/elementwise_linear_op.cu
+index cc49115bffc5..c1c45263f34c 100644
+--- a/caffe2/operators/elementwise_linear_op.cu
++++ b/caffe2/operators/elementwise_linear_op.cu
+@@ -5,6 +5,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/operator_fallback_gpu.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_mul_op.cu b/caffe2/operators/elementwise_mul_op.cu
+index bdbf760cf95b..88c3da00edc3 100644
+--- a/caffe2/operators/elementwise_mul_op.cu
++++ b/caffe2/operators/elementwise_mul_op.cu
+@@ -3,8 +3,8 @@
+ #include <algorithm>
+ #include <functional>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/elementwise_ops_utils.h"
+diff --git a/caffe2/operators/elementwise_ops.cu b/caffe2/operators/elementwise_ops.cu
+index c9ced33cf806..1ac0426d2ca7 100644
+--- a/caffe2/operators/elementwise_ops.cu
++++ b/caffe2/operators/elementwise_ops.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/elementwise_ops.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_load.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+diff --git a/caffe2/operators/find_op.cu b/caffe2/operators/find_op.cu
+index f8ff2bab1637..666df335ce42 100644
+--- a/caffe2/operators/find_op.cu
++++ b/caffe2/operators/find_op.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/find_op.h"
+diff --git a/caffe2/operators/generate_proposals_op.cu b/caffe2/operators/generate_proposals_op.cu
+index cab0ad3d0b88..84906a8e8182 100644
+--- a/caffe2/operators/generate_proposals_op.cu
++++ b/caffe2/operators/generate_proposals_op.cu
+@@ -1,4 +1,4 @@
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+ #include "caffe2/core/context.h"
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/generate_proposals_op.h"
+diff --git a/caffe2/operators/normalize_ops.cu b/caffe2/operators/normalize_ops.cu
+index 26df05308d88..468175df985f 100644
+--- a/caffe2/operators/normalize_ops.cu
++++ b/caffe2/operators/normalize_ops.cu
+@@ -1,5 +1,6 @@
+ #include <algorithm>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/one_hot_ops.cu b/caffe2/operators/one_hot_ops.cu
+index e521b3dd09df..86f82f78bb82 100644
+--- a/caffe2/operators/one_hot_ops.cu
++++ b/caffe2/operators/one_hot_ops.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/pack_segments.cu b/caffe2/operators/pack_segments.cu
+index 7475100fd368..b9ed413d1e7b 100644
+--- a/caffe2/operators/pack_segments.cu
++++ b/caffe2/operators/pack_segments.cu
+@@ -1,4 +1,4 @@
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/pack_segments.h"
+
+diff --git a/caffe2/operators/prelu_op.cu b/caffe2/operators/prelu_op.cu
+index 745a393f075f..d29882086754 100644
+--- a/caffe2/operators/prelu_op.cu
++++ b/caffe2/operators/prelu_op.cu
+@@ -1,6 +1,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/prelu_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/reduce_front_back_max_ops.cu b/caffe2/operators/reduce_front_back_max_ops.cu
+index ae91f8a6da72..ba62b2eff671 100644
+--- a/caffe2/operators/reduce_front_back_max_ops.cu
++++ b/caffe2/operators/reduce_front_back_max_ops.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_max_ops.h"
+diff --git a/caffe2/operators/reduce_front_back_sum_mean_ops.cu b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+index 476596f08425..586c20fe8d8e 100644
+--- a/caffe2/operators/reduce_front_back_sum_mean_ops.cu
++++ b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_sum_mean_ops.h"
+diff --git a/caffe2/operators/reduction_ops.cu b/caffe2/operators/reduction_ops.cu
+index ba55a66de588..0d94fab22a7f 100644
+--- a/caffe2/operators/reduction_ops.cu
++++ b/caffe2/operators/reduction_ops.cu
+@@ -2,7 +2,7 @@
+ #include "caffe2/operators/reduction_ops.h"
+ #include "caffe2/utils/conversions.h"
+
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/rmac_regions_op.cu b/caffe2/operators/rmac_regions_op.cu
+index 76c4d012d71a..39cc5fbc988d 100644
+--- a/caffe2/operators/rmac_regions_op.cu
++++ b/caffe2/operators/rmac_regions_op.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+@@ -10,6 +11,9 @@
+ #if defined(USE_ROCM)
+ namespace rocprim {
+ #else
++#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++namespace at_cuda_detail {
++#endif
+ namespace cub {
+ #endif
+
+@@ -22,6 +26,9 @@ inline __host__ __device__ bool operator<(
+ }
+
+ } // namespace cub
++#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++} // namespace at_cuda_detail
++#endif
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/segment_reduction_op_gpu.cuh b/caffe2/operators/segment_reduction_op_gpu.cuh
+index 8d51196ee138..447617c6e9de 100644
+--- a/caffe2/operators/segment_reduction_op_gpu.cuh
++++ b/caffe2/operators/segment_reduction_op_gpu.cuh
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+ #include <cub/device/device_scan.cuh>
+diff --git a/caffe2/operators/sequence_ops.cu b/caffe2/operators/sequence_ops.cu
+index cc34effd3f22..e66d491f85e6 100644
+--- a/caffe2/operators/sequence_ops.cu
++++ b/caffe2/operators/sequence_ops.cu
+@@ -1,6 +1,6 @@
+ #include <algorithm>
+
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/sequence_ops.h"
+diff --git a/caffe2/operators/softmax_ops.cu b/caffe2/operators/softmax_ops.cu
+index 51c0cbc2bf6a..c01fcf3e0a48 100644
+--- a/caffe2/operators/softmax_ops.cu
++++ b/caffe2/operators/softmax_ops.cu
+@@ -1,4 +1,5 @@
+ #include <cfloat>
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/spatial_batch_norm_op_impl.cuh b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+index edc076c7d718..f9b9fb58adc8 100644
+--- a/caffe2/operators/spatial_batch_norm_op_impl.cuh
++++ b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+@@ -5,8 +5,8 @@
+
+ #include <limits>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_fused_op_gpu.cu b/caffe2/sgd/adagrad_fused_op_gpu.cu
+index 2c2ad2cf76ae..396da5195125 100644
+--- a/caffe2/sgd/adagrad_fused_op_gpu.cu
++++ b/caffe2/sgd/adagrad_fused_op_gpu.cu
+@@ -2,6 +2,7 @@
+ #include <c10/core/GeneratorImpl.h>
+ #include <algorithm>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/device/device_radix_sort.cuh>
+ #include "caffe2/sgd/adagrad_fused_op_gpu.cuh"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_op_gpu.cu b/caffe2/sgd/adagrad_op_gpu.cu
+index 8abb3376ca87..a6fa842ddc80 100644
+--- a/caffe2/sgd/adagrad_op_gpu.cu
++++ b/caffe2/sgd/adagrad_op_gpu.cu
+@@ -1,5 +1,6 @@
+ #include <algorithm>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/sgd/adam_op_gpu.cu b/caffe2/sgd/adam_op_gpu.cu
+index 42ab975faacb..4b59836b6a68 100644
+--- a/caffe2/sgd/adam_op_gpu.cu
++++ b/caffe2/sgd/adam_op_gpu.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/utils/math/reduce.cu b/caffe2/utils/math/reduce.cu
+index fc3e476b288b..20919334da50 100644
+--- a/caffe2/utils/math/reduce.cu
++++ b/caffe2/utils/math/reduce.cu
+@@ -6,8 +6,8 @@
+ #include <numeric>
+ #include <vector>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include <thrust/execution_policy.h>
+ #include <thrust/reduce.h>
+diff --git a/caffe2/utils/math/reduce.cuh b/caffe2/utils/math/reduce.cuh
+index 0c43ad45a379..39ad553eec76 100644
+--- a/caffe2/utils/math/reduce.cuh
++++ b/caffe2/utils/math/reduce.cuh
+@@ -1,8 +1,8 @@
+ #ifndef CAFFE2_UTILS_MATH_REDUCE_CUH_
+ #define CAFFE2_UTILS_MATH_REDUCE_CUH_
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+
+diff --git a/caffe2/utils/math_gpu.cu b/caffe2/utils/math_gpu.cu
+index a37d4b744d73..b0a44fed34fb 100644
+--- a/caffe2/utils/math_gpu.cu
++++ b/caffe2/utils/math_gpu.cu
+@@ -7,8 +7,8 @@
+ #include <numeric>
+ #include <vector>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include <thrust/host_vector.h>
+ #include <thrust/device_vector.h>
+diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
+index ca560288a41a..f127d8f2e5b1 100644
+--- a/cmake/Dependencies.cmake
++++ b/cmake/Dependencies.cmake
+@@ -1622,6 +1622,11 @@ if(NOT INTERN_BUILD_MOBILE)
+ list(APPEND CUDA_NVCC_FLAGS "-Xcompiler" "-fPIC")
+ endif()
+
++ # include cub in a safe manner, see:
++ # https://github.com/pytorch/pytorch/pull/55292
++ # https://github.com/NVIDIA/cub/releases/tag/1.14.0
++ list(APPEND CUDA_NVCC_FLAGS "-DCUB_WRAPPED_NAMESPACE=at_cuda_detail")
++
+ if(CUDA_HAS_FP16 OR NOT ${CUDA_VERSION} LESS 7.5)
+ message(STATUS "Found CUDA with FP16 support, compiling with torch.cuda.HalfTensor")
+ list(APPEND CUDA_NVCC_FLAGS "-DCUDA_HAS_FP16=1" "-D__CUDA_NO_HALF_OPERATORS__" "-D__CUDA_NO_HALF_CONVERSIONS__"
+
+From f0a6afd7f0a5c6210289869282c77a680d7dfa2e Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Wed, 6 Oct 2021 12:16:48 -0700
+Subject: [PATCH 02/30] fix
+
+---
+ .github/workflows/lint.yml | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/.github/workflows/lint.yml b/.github/workflows/lint.yml
+index 4c469878e63f..af472e00140a 100644
+--- a/.github/workflows/lint.yml
++++ b/.github/workflows/lint.yml
+@@ -97,7 +97,7 @@ jobs:
+ - name: Ensure no direct cub include
+ if: always()
+ run: |
+- (! git --no-pager grep -I -no $'#include <cub/' -- ./aten ':(exclude)aten/src/ATen/cuda/cub.cuh' || (echo "The above files have direct cub include; please include ATen/cuda/cub.cuh instead and wrap your cub calls in at::native namespace if necessary"; false))
++ (! git --no-pager grep -I -no $'#include <cub/' -- ./aten ':(exclude)aten/src/ATen/cuda/cub.cuh' ':(exclude)aten/src/ATen/cuda/cub_definitions.cuh' || (echo "The above files have direct cub include; please include ATen/cuda/cub.cuh instead and wrap your cub calls in at::native namespace if necessary"; false))
+ - name: Ensure no raw cuda api calls
+ if: always()
+ run: |
+
+From 8edc6e961f1911f0fe04cb57af13b251f66c9153 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Wed, 6 Oct 2021 12:26:19 -0700
+Subject: [PATCH 03/30] save
+
+---
+ .github/workflows/lint.yml | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/.github/workflows/lint.yml b/.github/workflows/lint.yml
+index af472e00140a..0b2abe7a7483 100644
+--- a/.github/workflows/lint.yml
++++ b/.github/workflows/lint.yml
+@@ -97,7 +97,7 @@ jobs:
+ - name: Ensure no direct cub include
+ if: always()
+ run: |
+- (! git --no-pager grep -I -no $'#include <cub/' -- ./aten ':(exclude)aten/src/ATen/cuda/cub.cuh' ':(exclude)aten/src/ATen/cuda/cub_definitions.cuh' || (echo "The above files have direct cub include; please include ATen/cuda/cub.cuh instead and wrap your cub calls in at::native namespace if necessary"; false))
++ (! git --no-pager grep -I -no $'#include <cub/' -- ./aten ':(exclude)aten/src/ATen/cuda/cub*.cuh' || (echo "The above files have direct cub include; please include ATen/cuda/cub.cuh instead and wrap your cub calls in at::native namespace if necessary"; false))
+ - name: Ensure no raw cuda api calls
+ if: always()
+ run: |
+
+From 0d50954a3747ed0f2c7cbfcd58f17a2a81d5929c Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Wed, 6 Oct 2021 12:32:16 -0700
+Subject: [PATCH 04/30] fix
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index a7694fda4168..802f67f22a0f 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -14,4 +14,4 @@
+ #define CUB_SUPPORTS_WRAPPED_NAMESPACE() true
+ #elif
+ #define CUB_SUPPORTS_WRAPPED_NAMESPACE() false
+-#endif
+\ No newline at end of file
++#endif
+
+From 2ac5cf508d694995eee4710d9a6499bb3135d324 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Wed, 6 Oct 2021 14:04:40 -0700
+Subject: [PATCH 05/30] fix
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 4 ++--
+ 1 file changed, 2 insertions(+), 2 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 802f67f22a0f..1f5628fa9d3a 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -4,7 +4,7 @@
+ // https://github.com/NVIDIA/cub/pull/306
+ #if CUB_VERSION >= 101300
+ #define CUB_SUPPORTS_NV_BFLOAT16() true
+-#elif
++#else
+ #define CUB_SUPPORTS_NV_BFLOAT16() false
+ #endif
+
+@@ -12,6 +12,6 @@
+ // https://github.com/NVIDIA/cub/pull/326
+ #if CUB_VERSION >= 101400
+ #define CUB_SUPPORTS_WRAPPED_NAMESPACE() true
+-#elif
++#else
+ #define CUB_SUPPORTS_WRAPPED_NAMESPACE() false
+ #endif
+
+From 815f5a5981919f61aaab2d65597826ddf0495ac9 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Wed, 6 Oct 2021 14:06:32 -0700
+Subject: [PATCH 06/30] fix rocm
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 6 ++++++
+ 1 file changed, 6 insertions(+)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 1f5628fa9d3a..7da3ab716e41 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -1,4 +1,10 @@
++#pragma once
++
++#if !defined(USE_ROCM)
+ #include <cub/version.cuh>
++#else
++#define CUB_VERSION 0
++#endif
+
+ // cub sort support for __nv_bfloat16 is added to cub 1.13 in:
+ // https://github.com/NVIDIA/cub/pull/306
+
+From 02e9ca26bdc658edd1d960dc5d31e277b9b9afc1 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Mon, 11 Oct 2021 10:06:27 -0700
+Subject: [PATCH 07/30] revert caffe2 change
+
+---
+ caffe2/core/context_gpu.cu | 1 -
+ caffe2/operators/accuracy_op.cu | 1 -
+ caffe2/operators/affine_channel_op.cu | 1 -
+ caffe2/operators/arg_ops.cu | 2 +-
+ caffe2/operators/batch_moments_op.cu | 1 -
+ caffe2/operators/batch_sparse_to_dense_op.cu | 1 -
+ caffe2/operators/boolean_mask_ops.cu | 2 +-
+ caffe2/operators/cross_entropy_op.cu | 1 -
+ caffe2/operators/distance_op.cu | 1 -
+ caffe2/operators/elementwise_div_op.cu | 2 +-
+ caffe2/operators/elementwise_linear_op.cu | 1 -
+ caffe2/operators/elementwise_mul_op.cu | 2 +-
+ caffe2/operators/elementwise_ops.cu | 1 -
+ caffe2/operators/find_op.cu | 1 -
+ caffe2/operators/generate_proposals_op.cu | 2 +-
+ caffe2/operators/normalize_ops.cu | 1 -
+ caffe2/operators/one_hot_ops.cu | 1 -
+ caffe2/operators/pack_segments.cu | 2 +-
+ caffe2/operators/prelu_op.cu | 1 -
+ caffe2/operators/reduce_front_back_max_ops.cu | 1 -
+ caffe2/operators/reduce_front_back_sum_mean_ops.cu | 1 -
+ caffe2/operators/reduction_ops.cu | 2 +-
+ caffe2/operators/rmac_regions_op.cu | 7 -------
+ caffe2/operators/segment_reduction_op_gpu.cuh | 1 -
+ caffe2/operators/sequence_ops.cu | 2 +-
+ caffe2/operators/softmax_ops.cu | 1 -
+ caffe2/operators/spatial_batch_norm_op_impl.cuh | 2 +-
+ caffe2/sgd/adagrad_fused_op_gpu.cu | 1 -
+ caffe2/sgd/adagrad_op_gpu.cu | 1 -
+ caffe2/sgd/adam_op_gpu.cu | 1 -
+ caffe2/utils/math/reduce.cu | 2 +-
+ caffe2/utils/math/reduce.cuh | 2 +-
+ caffe2/utils/math_gpu.cu | 2 +-
+ 33 files changed, 12 insertions(+), 39 deletions(-)
+
+diff --git a/caffe2/core/context_gpu.cu b/caffe2/core/context_gpu.cu
+index 475ed61ab4f7..c2b89945ada9 100644
+--- a/caffe2/core/context_gpu.cu
++++ b/caffe2/core/context_gpu.cu
+@@ -4,7 +4,6 @@
+ #include <string>
+ #include <unordered_map>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <ATen/Context.h>
+ #include <c10/cuda/CUDAFunctions.h>
+ #include <c10/cuda/CUDACachingAllocator.h>
+diff --git a/caffe2/operators/accuracy_op.cu b/caffe2/operators/accuracy_op.cu
+index 7ad2b09c238a..f06663d71a90 100644
+--- a/caffe2/operators/accuracy_op.cu
++++ b/caffe2/operators/accuracy_op.cu
+@@ -3,7 +3,6 @@
+ #include "caffe2/utils/GpuAtomics.cuh"
+ #include "caffe2/utils/math.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/affine_channel_op.cu b/caffe2/operators/affine_channel_op.cu
+index f3d9e22c7e8c..adf4ac55c0fc 100644
+--- a/caffe2/operators/affine_channel_op.cu
++++ b/caffe2/operators/affine_channel_op.cu
+@@ -1,6 +1,5 @@
+ #include "caffe2/operators/affine_channel_op.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/arg_ops.cu b/caffe2/operators/arg_ops.cu
+index fbefe0774376..7e90d25b836b 100644
+--- a/caffe2/operators/arg_ops.cu
++++ b/caffe2/operators/arg_ops.cu
+@@ -2,8 +2,8 @@
+
+ #include <limits>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_moments_op.cu b/caffe2/operators/batch_moments_op.cu
+index 65c43200e5bd..4b693b5c04e2 100644
+--- a/caffe2/operators/batch_moments_op.cu
++++ b/caffe2/operators/batch_moments_op.cu
+@@ -1,6 +1,5 @@
+ #include "caffe2/operators/batch_moments_op.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_sparse_to_dense_op.cu b/caffe2/operators/batch_sparse_to_dense_op.cu
+index 2cb09deb8668..aea2035a5d3d 100644
+--- a/caffe2/operators/batch_sparse_to_dense_op.cu
++++ b/caffe2/operators/batch_sparse_to_dense_op.cu
+@@ -1,6 +1,5 @@
+ #include "caffe2/operators/batch_sparse_to_dense_op.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/device/device_scan.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/boolean_mask_ops.cu b/caffe2/operators/boolean_mask_ops.cu
+index c87688f51d64..214b7c13ba3c 100644
+--- a/caffe2/operators/boolean_mask_ops.cu
++++ b/caffe2/operators/boolean_mask_ops.cu
+@@ -3,7 +3,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/boolean_mask_ops.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/cross_entropy_op.cu b/caffe2/operators/cross_entropy_op.cu
+index 95f3ffddbf1f..380e80399fc3 100644
+--- a/caffe2/operators/cross_entropy_op.cu
++++ b/caffe2/operators/cross_entropy_op.cu
+@@ -1,5 +1,4 @@
+ #include <assert.h>
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/distance_op.cu b/caffe2/operators/distance_op.cu
+index d94691d5a9d9..3a8bb337d541 100644
+--- a/caffe2/operators/distance_op.cu
++++ b/caffe2/operators/distance_op.cu
+@@ -4,7 +4,6 @@
+ #include "caffe2/operators/distance_op.h"
+ #include "caffe2/utils/conversions.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_div_op.cu b/caffe2/operators/elementwise_div_op.cu
+index ca9682326324..42b103a0f110 100644
+--- a/caffe2/operators/elementwise_div_op.cu
++++ b/caffe2/operators/elementwise_div_op.cu
+@@ -3,8 +3,8 @@
+ #include <algorithm>
+ #include <functional>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/elementwise_ops_utils.h"
+diff --git a/caffe2/operators/elementwise_linear_op.cu b/caffe2/operators/elementwise_linear_op.cu
+index c1c45263f34c..cc49115bffc5 100644
+--- a/caffe2/operators/elementwise_linear_op.cu
++++ b/caffe2/operators/elementwise_linear_op.cu
+@@ -5,7 +5,6 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/operator_fallback_gpu.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_mul_op.cu b/caffe2/operators/elementwise_mul_op.cu
+index 88c3da00edc3..bdbf760cf95b 100644
+--- a/caffe2/operators/elementwise_mul_op.cu
++++ b/caffe2/operators/elementwise_mul_op.cu
+@@ -3,8 +3,8 @@
+ #include <algorithm>
+ #include <functional>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/elementwise_ops_utils.h"
+diff --git a/caffe2/operators/elementwise_ops.cu b/caffe2/operators/elementwise_ops.cu
+index 1ac0426d2ca7..c9ced33cf806 100644
+--- a/caffe2/operators/elementwise_ops.cu
++++ b/caffe2/operators/elementwise_ops.cu
+@@ -1,6 +1,5 @@
+ #include "caffe2/operators/elementwise_ops.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_load.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+diff --git a/caffe2/operators/find_op.cu b/caffe2/operators/find_op.cu
+index 666df335ce42..f8ff2bab1637 100644
+--- a/caffe2/operators/find_op.cu
++++ b/caffe2/operators/find_op.cu
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/find_op.h"
+diff --git a/caffe2/operators/generate_proposals_op.cu b/caffe2/operators/generate_proposals_op.cu
+index 84906a8e8182..cab0ad3d0b88 100644
+--- a/caffe2/operators/generate_proposals_op.cu
++++ b/caffe2/operators/generate_proposals_op.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+ #include "caffe2/core/context.h"
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/generate_proposals_op.h"
+diff --git a/caffe2/operators/normalize_ops.cu b/caffe2/operators/normalize_ops.cu
+index 468175df985f..26df05308d88 100644
+--- a/caffe2/operators/normalize_ops.cu
++++ b/caffe2/operators/normalize_ops.cu
+@@ -1,6 +1,5 @@
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/one_hot_ops.cu b/caffe2/operators/one_hot_ops.cu
+index 86f82f78bb82..e521b3dd09df 100644
+--- a/caffe2/operators/one_hot_ops.cu
++++ b/caffe2/operators/one_hot_ops.cu
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/pack_segments.cu b/caffe2/operators/pack_segments.cu
+index b9ed413d1e7b..7475100fd368 100644
+--- a/caffe2/operators/pack_segments.cu
++++ b/caffe2/operators/pack_segments.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/pack_segments.h"
+
+diff --git a/caffe2/operators/prelu_op.cu b/caffe2/operators/prelu_op.cu
+index d29882086754..745a393f075f 100644
+--- a/caffe2/operators/prelu_op.cu
++++ b/caffe2/operators/prelu_op.cu
+@@ -1,7 +1,6 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/prelu_op.h"
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/reduce_front_back_max_ops.cu b/caffe2/operators/reduce_front_back_max_ops.cu
+index ba62b2eff671..ae91f8a6da72 100644
+--- a/caffe2/operators/reduce_front_back_max_ops.cu
++++ b/caffe2/operators/reduce_front_back_max_ops.cu
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_max_ops.h"
+diff --git a/caffe2/operators/reduce_front_back_sum_mean_ops.cu b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+index 586c20fe8d8e..476596f08425 100644
+--- a/caffe2/operators/reduce_front_back_sum_mean_ops.cu
++++ b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_sum_mean_ops.h"
+diff --git a/caffe2/operators/reduction_ops.cu b/caffe2/operators/reduction_ops.cu
+index 0d94fab22a7f..ba55a66de588 100644
+--- a/caffe2/operators/reduction_ops.cu
++++ b/caffe2/operators/reduction_ops.cu
+@@ -2,7 +2,7 @@
+ #include "caffe2/operators/reduction_ops.h"
+ #include "caffe2/utils/conversions.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/rmac_regions_op.cu b/caffe2/operators/rmac_regions_op.cu
+index 39cc5fbc988d..76c4d012d71a 100644
+--- a/caffe2/operators/rmac_regions_op.cu
++++ b/caffe2/operators/rmac_regions_op.cu
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+@@ -11,9 +10,6 @@
+ #if defined(USE_ROCM)
+ namespace rocprim {
+ #else
+-#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
+-namespace at_cuda_detail {
+-#endif
+ namespace cub {
+ #endif
+
+@@ -26,9 +22,6 @@ inline __host__ __device__ bool operator<(
+ }
+
+ } // namespace cub
+-#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
+-} // namespace at_cuda_detail
+-#endif
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/segment_reduction_op_gpu.cuh b/caffe2/operators/segment_reduction_op_gpu.cuh
+index 447617c6e9de..8d51196ee138 100644
+--- a/caffe2/operators/segment_reduction_op_gpu.cuh
++++ b/caffe2/operators/segment_reduction_op_gpu.cuh
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+ #include <cub/device/device_scan.cuh>
+diff --git a/caffe2/operators/sequence_ops.cu b/caffe2/operators/sequence_ops.cu
+index e66d491f85e6..cc34effd3f22 100644
+--- a/caffe2/operators/sequence_ops.cu
++++ b/caffe2/operators/sequence_ops.cu
+@@ -1,6 +1,6 @@
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/sequence_ops.h"
+diff --git a/caffe2/operators/softmax_ops.cu b/caffe2/operators/softmax_ops.cu
+index c01fcf3e0a48..51c0cbc2bf6a 100644
+--- a/caffe2/operators/softmax_ops.cu
++++ b/caffe2/operators/softmax_ops.cu
+@@ -1,5 +1,4 @@
+ #include <cfloat>
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/spatial_batch_norm_op_impl.cuh b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+index f9b9fb58adc8..edc076c7d718 100644
+--- a/caffe2/operators/spatial_batch_norm_op_impl.cuh
++++ b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+@@ -5,8 +5,8 @@
+
+ #include <limits>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_fused_op_gpu.cu b/caffe2/sgd/adagrad_fused_op_gpu.cu
+index 396da5195125..2c2ad2cf76ae 100644
+--- a/caffe2/sgd/adagrad_fused_op_gpu.cu
++++ b/caffe2/sgd/adagrad_fused_op_gpu.cu
+@@ -2,7 +2,6 @@
+ #include <c10/core/GeneratorImpl.h>
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/device/device_radix_sort.cuh>
+ #include "caffe2/sgd/adagrad_fused_op_gpu.cuh"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_op_gpu.cu b/caffe2/sgd/adagrad_op_gpu.cu
+index a6fa842ddc80..8abb3376ca87 100644
+--- a/caffe2/sgd/adagrad_op_gpu.cu
++++ b/caffe2/sgd/adagrad_op_gpu.cu
+@@ -1,6 +1,5 @@
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/sgd/adam_op_gpu.cu b/caffe2/sgd/adam_op_gpu.cu
+index 4b59836b6a68..42ab975faacb 100644
+--- a/caffe2/sgd/adam_op_gpu.cu
++++ b/caffe2/sgd/adam_op_gpu.cu
+@@ -1,4 +1,3 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/utils/math/reduce.cu b/caffe2/utils/math/reduce.cu
+index 20919334da50..fc3e476b288b 100644
+--- a/caffe2/utils/math/reduce.cu
++++ b/caffe2/utils/math/reduce.cu
+@@ -6,8 +6,8 @@
+ #include <numeric>
+ #include <vector>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include <thrust/execution_policy.h>
+ #include <thrust/reduce.h>
+diff --git a/caffe2/utils/math/reduce.cuh b/caffe2/utils/math/reduce.cuh
+index 39ad553eec76..0c43ad45a379 100644
+--- a/caffe2/utils/math/reduce.cuh
++++ b/caffe2/utils/math/reduce.cuh
+@@ -1,8 +1,8 @@
+ #ifndef CAFFE2_UTILS_MATH_REDUCE_CUH_
+ #define CAFFE2_UTILS_MATH_REDUCE_CUH_
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+
+diff --git a/caffe2/utils/math_gpu.cu b/caffe2/utils/math_gpu.cu
+index b0a44fed34fb..a37d4b744d73 100644
+--- a/caffe2/utils/math_gpu.cu
++++ b/caffe2/utils/math_gpu.cu
+@@ -7,8 +7,8 @@
+ #include <numeric>
+ #include <vector>
+
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include <cub/cub.cuh>
+
+ #include <thrust/host_vector.h>
+ #include <thrust/device_vector.h>
+
+From a9ca6d97ff78a4a7a6a6dca90b21ad24666bf6fe Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Mon, 11 Oct 2021 10:27:36 -0700
+Subject: [PATCH 08/30] save
+
+---
+ aten/src/ATen/cuda/cub.cuh | 1 +
+ 1 file changed, 1 insertion(+)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 39938efc48be..23e90a81a16c 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -9,6 +9,7 @@
+
+ #if CUB_SUPPORTS_WRAPPED_NAMESPACE()
+
++#define CUB_WRAPPED_NAMESPACE at_cuda_detail
+ #include <cub/cub.cuh>
+
+ #else
+
+From 3695222a89a011cf58dfdb907515ffc72231f3af Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Mon, 11 Oct 2021 10:51:25 -0700
+Subject: [PATCH 09/30] fix
+
+---
+ aten/src/ATen/cuda/cub.cuh | 5 ++---
+ 1 file changed, 2 insertions(+), 3 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 23e90a81a16c..9014d4b7d2ec 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -67,12 +67,11 @@ template <> struct cub::NumericTraits<c10::BFloat16>: cub::BaseTraits<cub::FLOAT
+ }
+ #endif
+
++#if !defined(USE_ROCM)
+ namespace at { namespace native {
+ namespace cub = at_cuda_detail::cub;
+ }}
+-namespace caffe2 {
+-namespace cub = at_cuda_detail::cub;
+-}
++#endif
+
+ namespace at {
+ namespace cuda {
+
+From 4d94df2cf61598bb6078364a0f628cb0321cdff5 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 07:54:47 -0700
+Subject: [PATCH 10/30] Revert "revert caffe2 change"
+
+This reverts commit 02e9ca26bdc658edd1d960dc5d31e277b9b9afc1.
+---
+ caffe2/core/context_gpu.cu | 1 +
+ caffe2/operators/accuracy_op.cu | 1 +
+ caffe2/operators/affine_channel_op.cu | 1 +
+ caffe2/operators/arg_ops.cu | 2 +-
+ caffe2/operators/batch_moments_op.cu | 1 +
+ caffe2/operators/batch_sparse_to_dense_op.cu | 1 +
+ caffe2/operators/boolean_mask_ops.cu | 2 +-
+ caffe2/operators/cross_entropy_op.cu | 1 +
+ caffe2/operators/distance_op.cu | 1 +
+ caffe2/operators/elementwise_div_op.cu | 2 +-
+ caffe2/operators/elementwise_linear_op.cu | 1 +
+ caffe2/operators/elementwise_mul_op.cu | 2 +-
+ caffe2/operators/elementwise_ops.cu | 1 +
+ caffe2/operators/find_op.cu | 1 +
+ caffe2/operators/generate_proposals_op.cu | 2 +-
+ caffe2/operators/normalize_ops.cu | 1 +
+ caffe2/operators/one_hot_ops.cu | 1 +
+ caffe2/operators/pack_segments.cu | 2 +-
+ caffe2/operators/prelu_op.cu | 1 +
+ caffe2/operators/reduce_front_back_max_ops.cu | 1 +
+ caffe2/operators/reduce_front_back_sum_mean_ops.cu | 1 +
+ caffe2/operators/reduction_ops.cu | 2 +-
+ caffe2/operators/rmac_regions_op.cu | 7 +++++++
+ caffe2/operators/segment_reduction_op_gpu.cuh | 1 +
+ caffe2/operators/sequence_ops.cu | 2 +-
+ caffe2/operators/softmax_ops.cu | 1 +
+ caffe2/operators/spatial_batch_norm_op_impl.cuh | 2 +-
+ caffe2/sgd/adagrad_fused_op_gpu.cu | 1 +
+ caffe2/sgd/adagrad_op_gpu.cu | 1 +
+ caffe2/sgd/adam_op_gpu.cu | 1 +
+ caffe2/utils/math/reduce.cu | 2 +-
+ caffe2/utils/math/reduce.cuh | 2 +-
+ caffe2/utils/math_gpu.cu | 2 +-
+ 33 files changed, 39 insertions(+), 12 deletions(-)
+
+diff --git a/caffe2/core/context_gpu.cu b/caffe2/core/context_gpu.cu
+index c2b89945ada9..475ed61ab4f7 100644
+--- a/caffe2/core/context_gpu.cu
++++ b/caffe2/core/context_gpu.cu
+@@ -4,6 +4,7 @@
+ #include <string>
+ #include <unordered_map>
+
++#include <ATen/cuda/cub.cuh>
+ #include <ATen/Context.h>
+ #include <c10/cuda/CUDAFunctions.h>
+ #include <c10/cuda/CUDACachingAllocator.h>
+diff --git a/caffe2/operators/accuracy_op.cu b/caffe2/operators/accuracy_op.cu
+index f06663d71a90..7ad2b09c238a 100644
+--- a/caffe2/operators/accuracy_op.cu
++++ b/caffe2/operators/accuracy_op.cu
+@@ -3,6 +3,7 @@
+ #include "caffe2/utils/GpuAtomics.cuh"
+ #include "caffe2/utils/math.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/affine_channel_op.cu b/caffe2/operators/affine_channel_op.cu
+index adf4ac55c0fc..f3d9e22c7e8c 100644
+--- a/caffe2/operators/affine_channel_op.cu
++++ b/caffe2/operators/affine_channel_op.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/affine_channel_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/arg_ops.cu b/caffe2/operators/arg_ops.cu
+index 7e90d25b836b..fbefe0774376 100644
+--- a/caffe2/operators/arg_ops.cu
++++ b/caffe2/operators/arg_ops.cu
+@@ -2,8 +2,8 @@
+
+ #include <limits>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_moments_op.cu b/caffe2/operators/batch_moments_op.cu
+index 4b693b5c04e2..65c43200e5bd 100644
+--- a/caffe2/operators/batch_moments_op.cu
++++ b/caffe2/operators/batch_moments_op.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/batch_moments_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_sparse_to_dense_op.cu b/caffe2/operators/batch_sparse_to_dense_op.cu
+index aea2035a5d3d..2cb09deb8668 100644
+--- a/caffe2/operators/batch_sparse_to_dense_op.cu
++++ b/caffe2/operators/batch_sparse_to_dense_op.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/batch_sparse_to_dense_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/device/device_scan.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/boolean_mask_ops.cu b/caffe2/operators/boolean_mask_ops.cu
+index 214b7c13ba3c..c87688f51d64 100644
+--- a/caffe2/operators/boolean_mask_ops.cu
++++ b/caffe2/operators/boolean_mask_ops.cu
+@@ -3,7 +3,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/boolean_mask_ops.h"
+
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/cross_entropy_op.cu b/caffe2/operators/cross_entropy_op.cu
+index 380e80399fc3..95f3ffddbf1f 100644
+--- a/caffe2/operators/cross_entropy_op.cu
++++ b/caffe2/operators/cross_entropy_op.cu
+@@ -1,4 +1,5 @@
+ #include <assert.h>
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/distance_op.cu b/caffe2/operators/distance_op.cu
+index 3a8bb337d541..d94691d5a9d9 100644
+--- a/caffe2/operators/distance_op.cu
++++ b/caffe2/operators/distance_op.cu
+@@ -4,6 +4,7 @@
+ #include "caffe2/operators/distance_op.h"
+ #include "caffe2/utils/conversions.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_div_op.cu b/caffe2/operators/elementwise_div_op.cu
+index 42b103a0f110..ca9682326324 100644
+--- a/caffe2/operators/elementwise_div_op.cu
++++ b/caffe2/operators/elementwise_div_op.cu
+@@ -3,8 +3,8 @@
+ #include <algorithm>
+ #include <functional>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/elementwise_ops_utils.h"
+diff --git a/caffe2/operators/elementwise_linear_op.cu b/caffe2/operators/elementwise_linear_op.cu
+index cc49115bffc5..c1c45263f34c 100644
+--- a/caffe2/operators/elementwise_linear_op.cu
++++ b/caffe2/operators/elementwise_linear_op.cu
+@@ -5,6 +5,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/operator_fallback_gpu.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_mul_op.cu b/caffe2/operators/elementwise_mul_op.cu
+index bdbf760cf95b..88c3da00edc3 100644
+--- a/caffe2/operators/elementwise_mul_op.cu
++++ b/caffe2/operators/elementwise_mul_op.cu
+@@ -3,8 +3,8 @@
+ #include <algorithm>
+ #include <functional>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/elementwise_ops_utils.h"
+diff --git a/caffe2/operators/elementwise_ops.cu b/caffe2/operators/elementwise_ops.cu
+index c9ced33cf806..1ac0426d2ca7 100644
+--- a/caffe2/operators/elementwise_ops.cu
++++ b/caffe2/operators/elementwise_ops.cu
+@@ -1,5 +1,6 @@
+ #include "caffe2/operators/elementwise_ops.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_load.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+diff --git a/caffe2/operators/find_op.cu b/caffe2/operators/find_op.cu
+index f8ff2bab1637..666df335ce42 100644
+--- a/caffe2/operators/find_op.cu
++++ b/caffe2/operators/find_op.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/find_op.h"
+diff --git a/caffe2/operators/generate_proposals_op.cu b/caffe2/operators/generate_proposals_op.cu
+index cab0ad3d0b88..84906a8e8182 100644
+--- a/caffe2/operators/generate_proposals_op.cu
++++ b/caffe2/operators/generate_proposals_op.cu
+@@ -1,4 +1,4 @@
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+ #include "caffe2/core/context.h"
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/generate_proposals_op.h"
+diff --git a/caffe2/operators/normalize_ops.cu b/caffe2/operators/normalize_ops.cu
+index 26df05308d88..468175df985f 100644
+--- a/caffe2/operators/normalize_ops.cu
++++ b/caffe2/operators/normalize_ops.cu
+@@ -1,5 +1,6 @@
+ #include <algorithm>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/one_hot_ops.cu b/caffe2/operators/one_hot_ops.cu
+index e521b3dd09df..86f82f78bb82 100644
+--- a/caffe2/operators/one_hot_ops.cu
++++ b/caffe2/operators/one_hot_ops.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/pack_segments.cu b/caffe2/operators/pack_segments.cu
+index 7475100fd368..b9ed413d1e7b 100644
+--- a/caffe2/operators/pack_segments.cu
++++ b/caffe2/operators/pack_segments.cu
+@@ -1,4 +1,4 @@
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/pack_segments.h"
+
+diff --git a/caffe2/operators/prelu_op.cu b/caffe2/operators/prelu_op.cu
+index 745a393f075f..d29882086754 100644
+--- a/caffe2/operators/prelu_op.cu
++++ b/caffe2/operators/prelu_op.cu
+@@ -1,6 +1,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/prelu_op.h"
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/reduce_front_back_max_ops.cu b/caffe2/operators/reduce_front_back_max_ops.cu
+index ae91f8a6da72..ba62b2eff671 100644
+--- a/caffe2/operators/reduce_front_back_max_ops.cu
++++ b/caffe2/operators/reduce_front_back_max_ops.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_max_ops.h"
+diff --git a/caffe2/operators/reduce_front_back_sum_mean_ops.cu b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+index 476596f08425..586c20fe8d8e 100644
+--- a/caffe2/operators/reduce_front_back_sum_mean_ops.cu
++++ b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_sum_mean_ops.h"
+diff --git a/caffe2/operators/reduction_ops.cu b/caffe2/operators/reduction_ops.cu
+index ba55a66de588..0d94fab22a7f 100644
+--- a/caffe2/operators/reduction_ops.cu
++++ b/caffe2/operators/reduction_ops.cu
+@@ -2,7 +2,7 @@
+ #include "caffe2/operators/reduction_ops.h"
+ #include "caffe2/utils/conversions.h"
+
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/rmac_regions_op.cu b/caffe2/operators/rmac_regions_op.cu
+index 76c4d012d71a..39cc5fbc988d 100644
+--- a/caffe2/operators/rmac_regions_op.cu
++++ b/caffe2/operators/rmac_regions_op.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+@@ -10,6 +11,9 @@
+ #if defined(USE_ROCM)
+ namespace rocprim {
+ #else
++#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++namespace at_cuda_detail {
++#endif
+ namespace cub {
+ #endif
+
+@@ -22,6 +26,9 @@ inline __host__ __device__ bool operator<(
+ }
+
+ } // namespace cub
++#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++} // namespace at_cuda_detail
++#endif
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/segment_reduction_op_gpu.cuh b/caffe2/operators/segment_reduction_op_gpu.cuh
+index 8d51196ee138..447617c6e9de 100644
+--- a/caffe2/operators/segment_reduction_op_gpu.cuh
++++ b/caffe2/operators/segment_reduction_op_gpu.cuh
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+ #include <cub/device/device_scan.cuh>
+diff --git a/caffe2/operators/sequence_ops.cu b/caffe2/operators/sequence_ops.cu
+index cc34effd3f22..e66d491f85e6 100644
+--- a/caffe2/operators/sequence_ops.cu
++++ b/caffe2/operators/sequence_ops.cu
+@@ -1,6 +1,6 @@
+ #include <algorithm>
+
+-#include <cub/cub.cuh>
++#include <ATen/cuda/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/sequence_ops.h"
+diff --git a/caffe2/operators/softmax_ops.cu b/caffe2/operators/softmax_ops.cu
+index 51c0cbc2bf6a..c01fcf3e0a48 100644
+--- a/caffe2/operators/softmax_ops.cu
++++ b/caffe2/operators/softmax_ops.cu
+@@ -1,4 +1,5 @@
+ #include <cfloat>
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/spatial_batch_norm_op_impl.cuh b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+index edc076c7d718..f9b9fb58adc8 100644
+--- a/caffe2/operators/spatial_batch_norm_op_impl.cuh
++++ b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+@@ -5,8 +5,8 @@
+
+ #include <limits>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_fused_op_gpu.cu b/caffe2/sgd/adagrad_fused_op_gpu.cu
+index 2c2ad2cf76ae..396da5195125 100644
+--- a/caffe2/sgd/adagrad_fused_op_gpu.cu
++++ b/caffe2/sgd/adagrad_fused_op_gpu.cu
+@@ -2,6 +2,7 @@
+ #include <c10/core/GeneratorImpl.h>
+ #include <algorithm>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/device/device_radix_sort.cuh>
+ #include "caffe2/sgd/adagrad_fused_op_gpu.cuh"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_op_gpu.cu b/caffe2/sgd/adagrad_op_gpu.cu
+index 8abb3376ca87..a6fa842ddc80 100644
+--- a/caffe2/sgd/adagrad_op_gpu.cu
++++ b/caffe2/sgd/adagrad_op_gpu.cu
+@@ -1,5 +1,6 @@
+ #include <algorithm>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/sgd/adam_op_gpu.cu b/caffe2/sgd/adam_op_gpu.cu
+index 42ab975faacb..4b59836b6a68 100644
+--- a/caffe2/sgd/adam_op_gpu.cu
++++ b/caffe2/sgd/adam_op_gpu.cu
+@@ -1,3 +1,4 @@
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/utils/math/reduce.cu b/caffe2/utils/math/reduce.cu
+index fc3e476b288b..20919334da50 100644
+--- a/caffe2/utils/math/reduce.cu
++++ b/caffe2/utils/math/reduce.cu
+@@ -6,8 +6,8 @@
+ #include <numeric>
+ #include <vector>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include <thrust/execution_policy.h>
+ #include <thrust/reduce.h>
+diff --git a/caffe2/utils/math/reduce.cuh b/caffe2/utils/math/reduce.cuh
+index 0c43ad45a379..39ad553eec76 100644
+--- a/caffe2/utils/math/reduce.cuh
++++ b/caffe2/utils/math/reduce.cuh
+@@ -1,8 +1,8 @@
+ #ifndef CAFFE2_UTILS_MATH_REDUCE_CUH_
+ #define CAFFE2_UTILS_MATH_REDUCE_CUH_
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+
+diff --git a/caffe2/utils/math_gpu.cu b/caffe2/utils/math_gpu.cu
+index a37d4b744d73..b0a44fed34fb 100644
+--- a/caffe2/utils/math_gpu.cu
++++ b/caffe2/utils/math_gpu.cu
+@@ -7,8 +7,8 @@
+ #include <numeric>
+ #include <vector>
+
++#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
+-#include <cub/cub.cuh>
+
+ #include <thrust/host_vector.h>
+ #include <thrust/device_vector.h>
+
+From 359666307907780e946284445e609131652a4739 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 08:04:01 -0700
+Subject: [PATCH 11/30] save
+
+---
+ aten/src/ATen/cuda/cub.cuh | 4 +++-
+ aten/src/ATen/cuda/cub_definitions.cuh | 5 ++++-
+ cmake/Dependencies.cmake | 4 ++++
+ 3 files changed, 11 insertions(+), 2 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 9014d4b7d2ec..4f39bde4b3c8 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -9,7 +9,6 @@
+
+ #if CUB_SUPPORTS_WRAPPED_NAMESPACE()
+
+-#define CUB_WRAPPED_NAMESPACE at_cuda_detail
+ #include <cub/cub.cuh>
+
+ #else
+@@ -71,6 +70,9 @@ template <> struct cub::NumericTraits<c10::BFloat16>: cub::BaseTraits<cub::FLOAT
+ namespace at { namespace native {
+ namespace cub = at_cuda_detail::cub;
+ }}
++namespace caffew {
++namespace cub = at_cuda_detail::cub;
++}
+ #endif
+
+ namespace at {
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 7da3ab716e41..b4c82c9ae975 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -1,6 +1,7 @@
+ #pragma once
+
+ #if !defined(USE_ROCM)
++#include <cuda.h>
+ #include <cub/version.cuh>
+ #else
+ #define CUB_VERSION 0
+@@ -16,7 +17,9 @@
+
+ // cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.14 in:
+ // https://github.com/NVIDIA/cub/pull/326
+-#if CUB_VERSION >= 101400
++// CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
++// starting from CUDA 11.4
++#if defined(CUDA_VERSION) && CUDA_VERSION >= 11040
+ #define CUB_SUPPORTS_WRAPPED_NAMESPACE() true
+ #else
+ #define CUB_SUPPORTS_WRAPPED_NAMESPACE() false
+diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
+index 564fcebc0b08..bfee8f6dc78a 100644
+--- a/cmake/Dependencies.cmake
++++ b/cmake/Dependencies.cmake
+@@ -1618,6 +1618,10 @@ if(NOT INTERN_BUILD_MOBILE)
+ set(CMAKE_CXX_STANDARD 14)
+ endif()
+
++ if(NOT ${CUDA_VERSION} LESS 11.4)
++ string(APPEND CMAKE_CUDA_FLAGS " -DCUB_WRAPPED_NAMESPACE=at_cuda_detail")
++ endif()
++
+ if(CUDA_HAS_FP16 OR NOT ${CUDA_VERSION} LESS 7.5)
+ message(STATUS "Found CUDA with FP16 support, compiling with torch.cuda.HalfTensor")
+ string(APPEND CMAKE_CUDA_FLAGS " -DCUDA_HAS_FP16=1"
+
+From d4b3679d45cadbf9123a869cfefa001832bce04a Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 08:09:00 -0700
+Subject: [PATCH 12/30] save
+
+---
+ aten/src/ATen/cuda/cub.cuh | 2 +-
+ aten/src/ATen/cuda/cub_definitions.cuh | 4 ++--
+ 2 files changed, 3 insertions(+), 3 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 4f39bde4b3c8..2583b7e640ec 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -7,7 +7,7 @@
+
+ #include <ATen/cuda/cub_definitions.cuh>
+
+-#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++#if USE_GLOBAL_CUB_WRAPPED_NAMESPACE()
+
+ #include <cub/cub.cuh>
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index b4c82c9ae975..160d6f157396 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -20,7 +20,7 @@
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+ // starting from CUDA 11.4
+ #if defined(CUDA_VERSION) && CUDA_VERSION >= 11040
+-#define CUB_SUPPORTS_WRAPPED_NAMESPACE() true
++#define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+-#define CUB_SUPPORTS_WRAPPED_NAMESPACE() false
++#define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() false
+ #endif
+
+From 2b710dab756fceadf393d1c38cec91ef000e8fa8 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 08:23:39 -0700
+Subject: [PATCH 13/30] fix
+
+---
+ aten/src/ATen/cuda/cub.cuh | 6 +++---
+ 1 file changed, 3 insertions(+), 3 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 2583b7e640ec..042009f3d8b4 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -68,10 +68,10 @@ template <> struct cub::NumericTraits<c10::BFloat16>: cub::BaseTraits<cub::FLOAT
+
+ #if !defined(USE_ROCM)
+ namespace at { namespace native {
+-namespace cub = at_cuda_detail::cub;
++namespace cub = ::at_cuda_detail::cub;
+ }}
+-namespace caffew {
+-namespace cub = at_cuda_detail::cub;
++namespace caffe2 {
++namespace cub = ::at_cuda_detail::cub;
+ }
+ #endif
+
+
+From 34c57ca996c3249fe23e7037f98d1a9638187371 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 08:31:35 -0700
+Subject: [PATCH 14/30] fix
+
+---
+ caffe2/operators/rmac_regions_op.cu | 4 ++--
+ 1 file changed, 2 insertions(+), 2 deletions(-)
+
+diff --git a/caffe2/operators/rmac_regions_op.cu b/caffe2/operators/rmac_regions_op.cu
+index 39cc5fbc988d..2f6c230574b0 100644
+--- a/caffe2/operators/rmac_regions_op.cu
++++ b/caffe2/operators/rmac_regions_op.cu
+@@ -11,7 +11,7 @@
+ #if defined(USE_ROCM)
+ namespace rocprim {
+ #else
+-#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++#if USE_GLOBAL_CUB_WRAPPED_NAMESPACE()
+ namespace at_cuda_detail {
+ #endif
+ namespace cub {
+@@ -26,7 +26,7 @@ inline __host__ __device__ bool operator<(
+ }
+
+ } // namespace cub
+-#if CUB_SUPPORTS_WRAPPED_NAMESPACE()
++#if USE_GLOBAL_CUB_WRAPPED_NAMESPACE()
+ } // namespace at_cuda_detail
+ #endif
+
+
+From 3f6bce369e20fc75160c8151064e5cf3c4d871ce Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 08:56:27 -0700
+Subject: [PATCH 16/30] save
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 4 ++--
+ 1 file changed, 2 insertions(+), 2 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 160d6f157396..7746885d6428 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -15,11 +15,11 @@
+ #define CUB_SUPPORTS_NV_BFLOAT16() false
+ #endif
+
+-// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.14 in:
++// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.13.1 in:
+ // https://github.com/NVIDIA/cub/pull/326
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+ // starting from CUDA 11.4
+-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11040
++#if CUB_VERSION >= 101301
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() false
+
+From a2346ec6c5804085e49e81007ed70fd1aae0f333 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 09:04:04 -0700
+Subject: [PATCH 17/30] save
+
+---
+ aten/src/ATen/cuda/cub.cuh | 3 +--
+ 1 file changed, 1 insertion(+), 2 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 042009f3d8b4..17062586b0e3 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -17,9 +17,8 @@
+ #undef CUB_NS_PREFIX
+ #define CUB_NS_PREFIX namespace at_cuda_detail {
+ #define CUB_NS_POSTFIX }
++#define CUB_NS_QUALIFIER ::at_cuda_detail::cub
+ #include <cub/cub.cuh>
+-#undef CUB_NS_POSTFIX
+-#undef CUB_NS_PREFIX
+
+ #endif
+
+
+From add451ed3a579a794609bb26e82b31cec8527d03 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 09:06:53 -0700
+Subject: [PATCH 18/30] save
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 7746885d6428..fb43441148fc 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -19,7 +19,7 @@
+ // https://github.com/NVIDIA/cub/pull/326
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+ // starting from CUDA 11.4
+-#if CUB_VERSION >= 101301
++#if CUB_VERSION >= 101400
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() false
+
+From d1113be1236633978a51581ad9270069ddea5c30 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 09:07:12 -0700
+Subject: [PATCH 19/30] save
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index fb43441148fc..4b339b92c6dc 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -18,7 +18,7 @@
+ // cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.13.1 in:
+ // https://github.com/NVIDIA/cub/pull/326
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+-// starting from CUDA 11.4
++// starting from CUDA 11.5
+ #if CUB_VERSION >= 101400
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+
+From 59949c2c4945f7df89888bbbfb44060ea19e31ff Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 11:45:18 -0700
+Subject: [PATCH 20/30] save
+
+---
+ aten/src/ATen/cuda/cub.cuh | 8 ++++----
+ aten/src/ATen/cuda/cub_definitions.cuh | 8 ++++----
+ caffe2/core/context_gpu.cu | 2 +-
+ caffe2/operators/accuracy_op.cu | 2 +-
+ caffe2/operators/affine_channel_op.cu | 2 +-
+ caffe2/operators/arg_ops.cu | 2 +-
+ caffe2/operators/batch_moments_op.cu | 2 +-
+ caffe2/operators/batch_sparse_to_dense_op.cu | 2 +-
+ caffe2/operators/boolean_mask_ops.cu | 4 ++--
+ caffe2/operators/cross_entropy_op.cu | 2 +-
+ caffe2/operators/distance_op.cu | 2 +-
+ caffe2/operators/elementwise_div_op.cu | 2 +-
+ caffe2/operators/elementwise_linear_op.cu | 2 +-
+ caffe2/operators/elementwise_mul_op.cu | 2 +-
+ caffe2/operators/elementwise_ops.cu | 2 +-
+ caffe2/operators/find_op.cu | 2 +-
+ caffe2/operators/generate_proposals_op.cu | 3 ++-
+ caffe2/operators/normalize_ops.cu | 2 +-
+ caffe2/operators/one_hot_ops.cu | 2 +-
+ caffe2/operators/pack_segments.cu | 3 ++-
+ caffe2/operators/prelu_op.cu | 2 +-
+ caffe2/operators/reduce_front_back_max_ops.cu | 2 +-
+ caffe2/operators/reduce_front_back_sum_mean_ops.cu | 2 +-
+ caffe2/operators/reduction_ops.cu | 2 +-
+ caffe2/operators/rmac_regions_op.cu | 2 +-
+ caffe2/operators/segment_reduction_op_gpu.cuh | 2 +-
+ caffe2/operators/sequence_ops.cu | 3 ++-
+ caffe2/operators/softmax_ops.cu | 2 +-
+ caffe2/operators/spatial_batch_norm_op_impl.cuh | 2 +-
+ caffe2/sgd/adagrad_fused_op_gpu.cu | 2 +-
+ caffe2/sgd/adagrad_op_gpu.cu | 2 +-
+ caffe2/sgd/adam_op_gpu.cu | 2 +-
+ caffe2/utils/cub_namespace.cuh | 7 +++++++
+ caffe2/utils/math/reduce.cu | 3 +--
+ caffe2/utils/math/reduce.cuh | 2 +-
+ caffe2/utils/math_gpu.cu | 3 ++-
+ cmake/Dependencies.cmake | 2 +-
+ 37 files changed, 54 insertions(+), 44 deletions(-)
+ create mode 100644 caffe2/utils/cub_namespace.cuh
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index 17062586b0e3..de14455adb98 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -15,11 +15,14 @@
+
+ #undef CUB_NS_POSTFIX //undef to avoid redefinition warnings
+ #undef CUB_NS_PREFIX
++#undef CUB_NS_QUALIFIER
+ #define CUB_NS_PREFIX namespace at_cuda_detail {
+ #define CUB_NS_POSTFIX }
+ #define CUB_NS_QUALIFIER ::at_cuda_detail::cub
+ #include <cub/cub.cuh>
+-
++#undef CUB_NS_POSTFIX
++#undef CUB_NS_PREFIX
++#undef CUB_NS_QUALIFIER
+ #endif
+
+ #include <ATen/cuda/Exceptions.h>
+@@ -69,9 +72,6 @@ template <> struct cub::NumericTraits<c10::BFloat16>: cub::BaseTraits<cub::FLOAT
+ namespace at { namespace native {
+ namespace cub = ::at_cuda_detail::cub;
+ }}
+-namespace caffe2 {
+-namespace cub = ::at_cuda_detail::cub;
+-}
+ #endif
+
+ namespace at {
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 4b339b92c6dc..07cffe9d34e4 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -15,12 +15,12 @@
+ #define CUB_SUPPORTS_NV_BFLOAT16() false
+ #endif
+
+-// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.13.1 in:
++// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.14 in:
+ // https://github.com/NVIDIA/cub/pull/326
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+-// starting from CUDA 11.5
+-#if CUB_VERSION >= 101400
++// starting from CUDA 11.6
++#if defined(CUB_WRAPPED_NAMESPACE) || defined(THRUST_CUB_WRAPPED_NAMESPACE)
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() false
+-#endif
++#endif
+\ No newline at end of file
+diff --git a/caffe2/core/context_gpu.cu b/caffe2/core/context_gpu.cu
+index 475ed61ab4f7..9ba9f74d5376 100644
+--- a/caffe2/core/context_gpu.cu
++++ b/caffe2/core/context_gpu.cu
+@@ -4,7 +4,7 @@
+ #include <string>
+ #include <unordered_map>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <ATen/Context.h>
+ #include <c10/cuda/CUDAFunctions.h>
+ #include <c10/cuda/CUDACachingAllocator.h>
+diff --git a/caffe2/operators/accuracy_op.cu b/caffe2/operators/accuracy_op.cu
+index 7ad2b09c238a..29df54e752d3 100644
+--- a/caffe2/operators/accuracy_op.cu
++++ b/caffe2/operators/accuracy_op.cu
+@@ -3,7 +3,7 @@
+ #include "caffe2/utils/GpuAtomics.cuh"
+ #include "caffe2/utils/math.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/affine_channel_op.cu b/caffe2/operators/affine_channel_op.cu
+index f3d9e22c7e8c..efae0a3fc695 100644
+--- a/caffe2/operators/affine_channel_op.cu
++++ b/caffe2/operators/affine_channel_op.cu
+@@ -1,6 +1,6 @@
+ #include "caffe2/operators/affine_channel_op.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/arg_ops.cu b/caffe2/operators/arg_ops.cu
+index fbefe0774376..56deaa636356 100644
+--- a/caffe2/operators/arg_ops.cu
++++ b/caffe2/operators/arg_ops.cu
+@@ -2,7 +2,7 @@
+
+ #include <limits>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+diff --git a/caffe2/operators/batch_moments_op.cu b/caffe2/operators/batch_moments_op.cu
+index 65c43200e5bd..81359f6440b0 100644
+--- a/caffe2/operators/batch_moments_op.cu
++++ b/caffe2/operators/batch_moments_op.cu
+@@ -1,6 +1,6 @@
+ #include "caffe2/operators/batch_moments_op.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/batch_sparse_to_dense_op.cu b/caffe2/operators/batch_sparse_to_dense_op.cu
+index 2cb09deb8668..3e7ad8af9a5b 100644
+--- a/caffe2/operators/batch_sparse_to_dense_op.cu
++++ b/caffe2/operators/batch_sparse_to_dense_op.cu
+@@ -1,6 +1,6 @@
+ #include "caffe2/operators/batch_sparse_to_dense_op.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/device/device_scan.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/boolean_mask_ops.cu b/caffe2/operators/boolean_mask_ops.cu
+index c87688f51d64..501dd3b191c8 100644
+--- a/caffe2/operators/boolean_mask_ops.cu
++++ b/caffe2/operators/boolean_mask_ops.cu
+@@ -2,8 +2,8 @@
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/boolean_mask_ops.h"
+-
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/cross_entropy_op.cu b/caffe2/operators/cross_entropy_op.cu
+index 95f3ffddbf1f..15cb8a4f574a 100644
+--- a/caffe2/operators/cross_entropy_op.cu
++++ b/caffe2/operators/cross_entropy_op.cu
+@@ -1,5 +1,5 @@
+ #include <assert.h>
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/distance_op.cu b/caffe2/operators/distance_op.cu
+index d94691d5a9d9..a360166854ff 100644
+--- a/caffe2/operators/distance_op.cu
++++ b/caffe2/operators/distance_op.cu
+@@ -4,7 +4,7 @@
+ #include "caffe2/operators/distance_op.h"
+ #include "caffe2/utils/conversions.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_div_op.cu b/caffe2/operators/elementwise_div_op.cu
+index ca9682326324..33118a8f5e16 100644
+--- a/caffe2/operators/elementwise_div_op.cu
++++ b/caffe2/operators/elementwise_div_op.cu
+@@ -3,7 +3,7 @@
+ #include <algorithm>
+ #include <functional>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/elementwise_linear_op.cu b/caffe2/operators/elementwise_linear_op.cu
+index c1c45263f34c..8f749644b295 100644
+--- a/caffe2/operators/elementwise_linear_op.cu
++++ b/caffe2/operators/elementwise_linear_op.cu
+@@ -5,7 +5,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/operator_fallback_gpu.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/elementwise_mul_op.cu b/caffe2/operators/elementwise_mul_op.cu
+index 88c3da00edc3..1991b8b513af 100644
+--- a/caffe2/operators/elementwise_mul_op.cu
++++ b/caffe2/operators/elementwise_mul_op.cu
+@@ -3,7 +3,7 @@
+ #include <algorithm>
+ #include <functional>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/elementwise_ops.cu b/caffe2/operators/elementwise_ops.cu
+index 1ac0426d2ca7..932bd5dafda0 100644
+--- a/caffe2/operators/elementwise_ops.cu
++++ b/caffe2/operators/elementwise_ops.cu
+@@ -1,6 +1,6 @@
+ #include "caffe2/operators/elementwise_ops.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_load.cuh>
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+diff --git a/caffe2/operators/find_op.cu b/caffe2/operators/find_op.cu
+index 666df335ce42..20d42560b506 100644
+--- a/caffe2/operators/find_op.cu
++++ b/caffe2/operators/find_op.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/find_op.h"
+diff --git a/caffe2/operators/generate_proposals_op.cu b/caffe2/operators/generate_proposals_op.cu
+index 84906a8e8182..b63726651939 100644
+--- a/caffe2/operators/generate_proposals_op.cu
++++ b/caffe2/operators/generate_proposals_op.cu
+@@ -1,10 +1,11 @@
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+ #include "caffe2/core/context.h"
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/generate_proposals_op.h"
+ #include "caffe2/operators/generate_proposals_op_util_boxes.h" // BBOX_XFORM_CLIP_DEFAULT
+ #include "caffe2/operators/generate_proposals_op_util_nms.h"
+ #include "caffe2/operators/generate_proposals_op_util_nms_gpu.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ #if defined(USE_ROCM)
+ #include <cfloat>
+diff --git a/caffe2/operators/normalize_ops.cu b/caffe2/operators/normalize_ops.cu
+index 468175df985f..952c4a772fa5 100644
+--- a/caffe2/operators/normalize_ops.cu
++++ b/caffe2/operators/normalize_ops.cu
+@@ -1,6 +1,6 @@
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/one_hot_ops.cu b/caffe2/operators/one_hot_ops.cu
+index 86f82f78bb82..4b1e054b0806 100644
+--- a/caffe2/operators/one_hot_ops.cu
++++ b/caffe2/operators/one_hot_ops.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/pack_segments.cu b/caffe2/operators/pack_segments.cu
+index b9ed413d1e7b..372638abdd24 100644
+--- a/caffe2/operators/pack_segments.cu
++++ b/caffe2/operators/pack_segments.cu
+@@ -1,6 +1,7 @@
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/pack_segments.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/prelu_op.cu b/caffe2/operators/prelu_op.cu
+index d29882086754..6303b70b4a89 100644
+--- a/caffe2/operators/prelu_op.cu
++++ b/caffe2/operators/prelu_op.cu
+@@ -1,7 +1,7 @@
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/prelu_op.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ namespace caffe2 {
+diff --git a/caffe2/operators/reduce_front_back_max_ops.cu b/caffe2/operators/reduce_front_back_max_ops.cu
+index ba62b2eff671..d6bb862e4dbb 100644
+--- a/caffe2/operators/reduce_front_back_max_ops.cu
++++ b/caffe2/operators/reduce_front_back_max_ops.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_max_ops.h"
+diff --git a/caffe2/operators/reduce_front_back_sum_mean_ops.cu b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+index 586c20fe8d8e..2b5cb7110edf 100644
+--- a/caffe2/operators/reduce_front_back_sum_mean_ops.cu
++++ b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_sum_mean_ops.h"
+diff --git a/caffe2/operators/reduction_ops.cu b/caffe2/operators/reduction_ops.cu
+index 0d94fab22a7f..9649b85d015c 100644
+--- a/caffe2/operators/reduction_ops.cu
++++ b/caffe2/operators/reduction_ops.cu
+@@ -2,7 +2,7 @@
+ #include "caffe2/operators/reduction_ops.h"
+ #include "caffe2/utils/conversions.h"
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/rmac_regions_op.cu b/caffe2/operators/rmac_regions_op.cu
+index 2f6c230574b0..6b79e0c4c8b4 100644
+--- a/caffe2/operators/rmac_regions_op.cu
++++ b/caffe2/operators/rmac_regions_op.cu
+@@ -1,5 +1,5 @@
+-#include <ATen/cuda/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/rmac_regions_op.h"
+diff --git a/caffe2/operators/segment_reduction_op_gpu.cuh b/caffe2/operators/segment_reduction_op_gpu.cuh
+index 447617c6e9de..bb3f3be13c72 100644
+--- a/caffe2/operators/segment_reduction_op_gpu.cuh
++++ b/caffe2/operators/segment_reduction_op_gpu.cuh
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include <cub/device/device_reduce.cuh>
+ #include <cub/device/device_scan.cuh>
+diff --git a/caffe2/operators/sequence_ops.cu b/caffe2/operators/sequence_ops.cu
+index e66d491f85e6..2ceb5236ef72 100644
+--- a/caffe2/operators/sequence_ops.cu
++++ b/caffe2/operators/sequence_ops.cu
+@@ -1,6 +1,7 @@
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/sequence_ops.h"
+diff --git a/caffe2/operators/softmax_ops.cu b/caffe2/operators/softmax_ops.cu
+index c01fcf3e0a48..b0afac3332a6 100644
+--- a/caffe2/operators/softmax_ops.cu
++++ b/caffe2/operators/softmax_ops.cu
+@@ -1,5 +1,5 @@
+ #include <cfloat>
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/operators/spatial_batch_norm_op_impl.cuh b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+index f9b9fb58adc8..6fdb4c63f8ef 100644
+--- a/caffe2/operators/spatial_batch_norm_op_impl.cuh
++++ b/caffe2/operators/spatial_batch_norm_op_impl.cuh
+@@ -5,7 +5,7 @@
+
+ #include <limits>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/sgd/adagrad_fused_op_gpu.cu b/caffe2/sgd/adagrad_fused_op_gpu.cu
+index 396da5195125..63d0712e3970 100644
+--- a/caffe2/sgd/adagrad_fused_op_gpu.cu
++++ b/caffe2/sgd/adagrad_fused_op_gpu.cu
+@@ -2,7 +2,7 @@
+ #include <c10/core/GeneratorImpl.h>
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/device/device_radix_sort.cuh>
+ #include "caffe2/sgd/adagrad_fused_op_gpu.cuh"
+ #include "caffe2/utils/math.h"
+diff --git a/caffe2/sgd/adagrad_op_gpu.cu b/caffe2/sgd/adagrad_op_gpu.cu
+index a6fa842ddc80..0b7f499345be 100644
+--- a/caffe2/sgd/adagrad_op_gpu.cu
++++ b/caffe2/sgd/adagrad_op_gpu.cu
+@@ -1,6 +1,6 @@
+ #include <algorithm>
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/sgd/adam_op_gpu.cu b/caffe2/sgd/adam_op_gpu.cu
+index 4b59836b6a68..a93812fabbe8 100644
+--- a/caffe2/sgd/adam_op_gpu.cu
++++ b/caffe2/sgd/adam_op_gpu.cu
+@@ -1,4 +1,4 @@
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+diff --git a/caffe2/utils/cub_namespace.cuh b/caffe2/utils/cub_namespace.cuh
+new file mode 100644
+index 000000000000..c7a5db0dc013
+--- /dev/null
++++ b/caffe2/utils/cub_namespace.cuh
+@@ -0,0 +1,7 @@
++#include <ATen/cuda/cub_definitions.cuh>
++
++#if USE_GLOBAL_CUB_WRAPPED_NAMESPACE()
++namespace caffe2 {
++namespace cub = ::CUB_WRAPPED_NAMESPACE::cub;
++}
++#endif
+diff --git a/caffe2/utils/math/reduce.cu b/caffe2/utils/math/reduce.cu
+index 20919334da50..69a6469d8ed1 100644
+--- a/caffe2/utils/math/reduce.cu
++++ b/caffe2/utils/math/reduce.cu
+@@ -5,8 +5,7 @@
+ #include <limits>
+ #include <numeric>
+ #include <vector>
+-
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include <thrust/execution_policy.h>
+diff --git a/caffe2/utils/math/reduce.cuh b/caffe2/utils/math/reduce.cuh
+index 39ad553eec76..18bdca11b9de 100644
+--- a/caffe2/utils/math/reduce.cuh
++++ b/caffe2/utils/math/reduce.cuh
+@@ -1,7 +1,7 @@
+ #ifndef CAFFE2_UTILS_MATH_REDUCE_CUH_
+ #define CAFFE2_UTILS_MATH_REDUCE_CUH_
+
+-#include <ATen/cuda/cub.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/common_gpu.h"
+diff --git a/caffe2/utils/math_gpu.cu b/caffe2/utils/math_gpu.cu
+index b0a44fed34fb..54b0a9391c26 100644
+--- a/caffe2/utils/math_gpu.cu
++++ b/caffe2/utils/math_gpu.cu
+@@ -7,8 +7,9 @@
+ #include <numeric>
+ #include <vector>
+
+-#include <ATen/cuda/cub.cuh>
++#include <cub/cub.cuh>
+ #include <cub/block/block_reduce.cuh>
++#include "caffe2/utils/cub_namespace.cuh"
+
+ #include <thrust/host_vector.h>
+ #include <thrust/device_vector.h>
+diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
+index bfee8f6dc78a..a6a3946bac8e 100644
+--- a/cmake/Dependencies.cmake
++++ b/cmake/Dependencies.cmake
+@@ -1618,7 +1618,7 @@ if(NOT INTERN_BUILD_MOBILE)
+ set(CMAKE_CXX_STANDARD 14)
+ endif()
+
+- if(NOT ${CUDA_VERSION} LESS 11.4)
++ if(NOT ${CUDA_VERSION} LESS 11.6)
+ string(APPEND CMAKE_CUDA_FLAGS " -DCUB_WRAPPED_NAMESPACE=at_cuda_detail")
+ endif()
+
+
+From c22df90d62acfcb572b4b7bf55646fe338843a69 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 11:47:52 -0700
+Subject: [PATCH 21/30] save
+
+---
+ aten/src/ATen/cuda/cub.cuh | 1 +
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 2 files changed, 2 insertions(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index de14455adb98..fecf9e077a6f 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -23,6 +23,7 @@
+ #undef CUB_NS_POSTFIX
+ #undef CUB_NS_PREFIX
+ #undef CUB_NS_QUALIFIER
++
+ #endif
+
+ #include <ATen/cuda/Exceptions.h>
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 07cffe9d34e4..d639d28e66b6 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -23,4 +23,4 @@
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() false
+-#endif
+\ No newline at end of file
++#endif
+
+From 64adfaaaca8267cd6cf03ae7a524eee239530425 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 12:29:41 -0700
+Subject: [PATCH 22/30] save
+
+---
+ caffe2/core/context_gpu.cu | 2 +-
+ caffe2/operators/cross_entropy_op.cu | 2 +-
+ caffe2/operators/find_op.cu | 2 +-
+ caffe2/operators/normalize_ops.cu | 2 +-
+ caffe2/operators/one_hot_ops.cu | 2 +-
+ caffe2/operators/reduce_front_back_max_ops.cu | 2 +-
+ caffe2/operators/reduce_front_back_sum_mean_ops.cu | 2 +-
+ caffe2/operators/softmax_ops.cu | 2 +-
+ caffe2/sgd/adagrad_op_gpu.cu | 2 +-
+ caffe2/sgd/adam_op_gpu.cu | 2 +-
+ cmake/Dependencies.cmake | 2 +-
+ 11 files changed, 11 insertions(+), 11 deletions(-)
+
+diff --git a/caffe2/core/context_gpu.cu b/caffe2/core/context_gpu.cu
+index 9ba9f74d5376..6d537400913e 100644
+--- a/caffe2/core/context_gpu.cu
++++ b/caffe2/core/context_gpu.cu
+@@ -4,7 +4,6 @@
+ #include <string>
+ #include <unordered_map>
+
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <ATen/Context.h>
+ #include <c10/cuda/CUDAFunctions.h>
+ #include <c10/cuda/CUDACachingAllocator.h>
+@@ -22,6 +21,7 @@
+ #include "caffe2/core/logging.h"
+ #include "caffe2/core/tensor.h"
+ #include "caffe2/utils/string_utils.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ C10_DEFINE_string(
+ caffe2_cuda_memory_pool,
+diff --git a/caffe2/operators/cross_entropy_op.cu b/caffe2/operators/cross_entropy_op.cu
+index 15cb8a4f574a..c23f05f8e5c2 100644
+--- a/caffe2/operators/cross_entropy_op.cu
++++ b/caffe2/operators/cross_entropy_op.cu
+@@ -1,10 +1,10 @@
+ #include <assert.h>
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/cross_entropy_op.h"
+ #include "caffe2/operators/operator_fallback_gpu.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/find_op.cu b/caffe2/operators/find_op.cu
+index 20d42560b506..0418a71fbcda 100644
+--- a/caffe2/operators/find_op.cu
++++ b/caffe2/operators/find_op.cu
+@@ -1,7 +1,7 @@
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/find_op.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/normalize_ops.cu b/caffe2/operators/normalize_ops.cu
+index 952c4a772fa5..e4d1f34b754c 100644
+--- a/caffe2/operators/normalize_ops.cu
++++ b/caffe2/operators/normalize_ops.cu
+@@ -1,11 +1,11 @@
+ #include <algorithm>
+
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/normalize_l1_op.h"
+ #include "caffe2/operators/normalize_op.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/one_hot_ops.cu b/caffe2/operators/one_hot_ops.cu
+index 4b1e054b0806..87e8196765ef 100644
+--- a/caffe2/operators/one_hot_ops.cu
++++ b/caffe2/operators/one_hot_ops.cu
+@@ -1,8 +1,8 @@
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/one_hot_ops.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/reduce_front_back_max_ops.cu b/caffe2/operators/reduce_front_back_max_ops.cu
+index d6bb862e4dbb..2ea25de46009 100644
+--- a/caffe2/operators/reduce_front_back_max_ops.cu
++++ b/caffe2/operators/reduce_front_back_max_ops.cu
+@@ -1,7 +1,7 @@
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_max_ops.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ #if defined(USE_ROCM)
+ #include <cfloat>
+diff --git a/caffe2/operators/reduce_front_back_sum_mean_ops.cu b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+index 2b5cb7110edf..a7ad6dd50084 100644
+--- a/caffe2/operators/reduce_front_back_sum_mean_ops.cu
++++ b/caffe2/operators/reduce_front_back_sum_mean_ops.cu
+@@ -1,7 +1,7 @@
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/reduce_front_back_sum_mean_ops.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/operators/softmax_ops.cu b/caffe2/operators/softmax_ops.cu
+index b0afac3332a6..ebf0700c9ef0 100644
+--- a/caffe2/operators/softmax_ops.cu
++++ b/caffe2/operators/softmax_ops.cu
+@@ -1,11 +1,11 @@
+ #include <cfloat>
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/operators/softmax_op.h"
+ #include "caffe2/operators/softmax_with_loss_op.h"
+ #include "caffe2/operators/spatial_softmax_with_loss_op.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/sgd/adagrad_op_gpu.cu b/caffe2/sgd/adagrad_op_gpu.cu
+index 0b7f499345be..b80d29700c3f 100644
+--- a/caffe2/sgd/adagrad_op_gpu.cu
++++ b/caffe2/sgd/adagrad_op_gpu.cu
+@@ -1,10 +1,10 @@
+ #include <algorithm>
+
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/sgd/adagrad_op.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/caffe2/sgd/adam_op_gpu.cu b/caffe2/sgd/adam_op_gpu.cu
+index a93812fabbe8..6f9c3234204d 100644
+--- a/caffe2/sgd/adam_op_gpu.cu
++++ b/caffe2/sgd/adam_op_gpu.cu
+@@ -1,8 +1,8 @@
+-#include "caffe2/utils/cub_namespace.cuh"
+ #include <cub/block/block_reduce.cuh>
+ #include "caffe2/core/common_gpu.h"
+ #include "caffe2/core/context_gpu.h"
+ #include "caffe2/sgd/adam_op.h"
++#include "caffe2/utils/cub_namespace.cuh"
+
+ namespace caffe2 {
+
+diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
+index a6a3946bac8e..5074bab1e999 100644
+--- a/cmake/Dependencies.cmake
++++ b/cmake/Dependencies.cmake
+@@ -1618,7 +1618,7 @@ if(NOT INTERN_BUILD_MOBILE)
+ set(CMAKE_CXX_STANDARD 14)
+ endif()
+
+- if(NOT ${CUDA_VERSION} LESS 11.6)
++ if(NOT ${CUDA_VERSION} LESS 11.5)
+ string(APPEND CMAKE_CUDA_FLAGS " -DCUB_WRAPPED_NAMESPACE=at_cuda_detail")
+ endif()
+
+
+From 21abede594471e4f225def12796c1d7c73a96b20 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 12:46:32 -0700
+Subject: [PATCH 23/30] save
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 4 ++--
+ 1 file changed, 2 insertions(+), 2 deletions(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index d639d28e66b6..9828abdfc99a 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -15,10 +15,10 @@
+ #define CUB_SUPPORTS_NV_BFLOAT16() false
+ #endif
+
+-// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.14 in:
++// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.13.1 in:
+ // https://github.com/NVIDIA/cub/pull/326
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+-// starting from CUDA 11.6
++// starting from CUDA 11.5
+ #if defined(CUB_WRAPPED_NAMESPACE) || defined(THRUST_CUB_WRAPPED_NAMESPACE)
+ #define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
+ #else
+
+From c73b101377d5adb8f96eb31988ceef7c2e760839 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Tue, 12 Oct 2021 12:59:12 -0700
+Subject: [PATCH 24/30] comment
+
+---
+ aten/src/ATen/cuda/cub.cuh | 2 ++
+ cmake/Dependencies.cmake | 2 ++
+ 2 files changed, 4 insertions(+)
+
+diff --git a/aten/src/ATen/cuda/cub.cuh b/aten/src/ATen/cuda/cub.cuh
+index fecf9e077a6f..0532470d74ec 100644
+--- a/aten/src/ATen/cuda/cub.cuh
++++ b/aten/src/ATen/cuda/cub.cuh
+@@ -13,6 +13,8 @@
+
+ #else
+
++// include cub in a safe manner, see:
++// https://github.com/pytorch/pytorch/pull/55292
+ #undef CUB_NS_POSTFIX //undef to avoid redefinition warnings
+ #undef CUB_NS_PREFIX
+ #undef CUB_NS_QUALIFIER
+diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
+index 5074bab1e999..bde42bf1719f 100644
+--- a/cmake/Dependencies.cmake
++++ b/cmake/Dependencies.cmake
+@@ -1618,6 +1618,8 @@ if(NOT INTERN_BUILD_MOBILE)
+ set(CMAKE_CXX_STANDARD 14)
+ endif()
+
++ # use cub in a safe manner, see:
++ # https://github.com/pytorch/pytorch/pull/55292
+ if(NOT ${CUDA_VERSION} LESS 11.5)
+ string(APPEND CMAKE_CUDA_FLAGS " -DCUB_WRAPPED_NAMESPACE=at_cuda_detail")
+ endif()
+
+From 1cb29c9816e795767236854644a1ba2c28da6e0d Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Fri, 22 Oct 2021 10:42:32 -0700
+Subject: [PATCH 25/30] fix
+
+---
+ caffe2/utils/cub_namespace.cuh | 19 ++++++++++++++++++-
+ 1 file changed, 18 insertions(+), 1 deletion(-)
+
+diff --git a/caffe2/utils/cub_namespace.cuh b/caffe2/utils/cub_namespace.cuh
+index c7a5db0dc013..752f273128ee 100644
+--- a/caffe2/utils/cub_namespace.cuh
++++ b/caffe2/utils/cub_namespace.cuh
+@@ -1,4 +1,21 @@
+-#include <ATen/cuda/cub_definitions.cuh>
++#pragma once
++
++#if !defined(USE_ROCM)
++#include <cuda.h>
++#include <cub/version.cuh>
++#else
++#define CUB_VERSION 0
++#endif
++
++// cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.13.1 in:
++// https://github.com/NVIDIA/cub/pull/326
++// CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
++// starting from CUDA 11.5
++#if defined(CUB_WRAPPED_NAMESPACE) || defined(THRUST_CUB_WRAPPED_NAMESPACE)
++#define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() true
++#else
++#define USE_GLOBAL_CUB_WRAPPED_NAMESPACE() false
++#endif
+
+ #if USE_GLOBAL_CUB_WRAPPED_NAMESPACE()
+ namespace caffe2 {
+
+From 3ca182d3c6d1dcc45396c3f2f3169cc84496ca1f Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Sun, 24 Oct 2021 14:27:09 -0700
+Subject: [PATCH 26/30] fix some cuda 10 build
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 9828abdfc99a..5243e551f4a7 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -1,6 +1,6 @@
+ #pragma once
+
+-#if !defined(USE_ROCM)
++#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
+ #include <cuda.h>
+ #include <cub/version.cuh>
+ #else
+
+From 3a7c2bc6da6f2be060be879d0b2277b833ab7cc9 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Sun, 24 Oct 2021 14:28:53 -0700
+Subject: [PATCH 27/30] fix caffe2
+
+---
+ caffe2/utils/cub_namespace.cuh | 7 -------
+ 1 file changed, 7 deletions(-)
+
+diff --git a/caffe2/utils/cub_namespace.cuh b/caffe2/utils/cub_namespace.cuh
+index 752f273128ee..188a9936f9c6 100644
+--- a/caffe2/utils/cub_namespace.cuh
++++ b/caffe2/utils/cub_namespace.cuh
+@@ -1,12 +1,5 @@
+ #pragma once
+
+-#if !defined(USE_ROCM)
+-#include <cuda.h>
+-#include <cub/version.cuh>
+-#else
+-#define CUB_VERSION 0
+-#endif
+-
+ // cub sort support for CUB_WRAPPED_NAMESPACE is added to cub 1.13.1 in:
+ // https://github.com/NVIDIA/cub/pull/326
+ // CUB_WRAPPED_NAMESPACE is defined globally in cmake/Dependencies.cmake
+
+From 960606e0dc1da7c16b6cb0ae11b04bace95816ac Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Sun, 24 Oct 2021 14:50:13 -0700
+Subject: [PATCH 28/30] fix
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 5 ++++-
+ 1 file changed, 4 insertions(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 5243e551f4a7..b921af480655 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -1,7 +1,10 @@
+ #pragma once
+
+-#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
++#if !define(USE_ROCM)
+ #include <cuda.h>
++#endif
++
++#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
+ #include <cub/version.cuh>
+ #else
+ #define CUB_VERSION 0
+
+From 48f128f55396622bfb952b474888813334fae303 Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Sun, 24 Oct 2021 14:51:12 -0700
+Subject: [PATCH 29/30] save
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index b921af480655..8c99f4951c6c 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -1,7 +1,7 @@
+ #pragma once
+
+ #if !define(USE_ROCM)
+-#include <cuda.h>
++#include <cuda.h> // for CUDA_VERSION
+ #endif
+
+ #if defined(CUDA_VERSION) && CUDA_VERSION >= 11000
+
+From b2023a036afe9ddff4f710953ad6ba55f8396ecf Mon Sep 17 00:00:00 2001
+From: Xiang Gao <qasdfgtyuiop at gmail.com>
+Date: Sun, 24 Oct 2021 14:51:59 -0700
+Subject: [PATCH 30/30] fixed
+
+---
+ aten/src/ATen/cuda/cub_definitions.cuh | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/aten/src/ATen/cuda/cub_definitions.cuh b/aten/src/ATen/cuda/cub_definitions.cuh
+index 8c99f4951c6c..61119fc17458 100644
+--- a/aten/src/ATen/cuda/cub_definitions.cuh
++++ b/aten/src/ATen/cuda/cub_definitions.cuh
+@@ -1,6 +1,6 @@
+ #pragma once
+
+-#if !define(USE_ROCM)
++#if !defined(USE_ROCM)
+ #include <cuda.h> // for CUDA_VERSION
+ #endif
+
Modified: PKGBUILD
===================================================================
--- PKGBUILD 2021-10-31 22:13:18 UTC (rev 1035134)
+++ PKGBUILD 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -2,12 +2,13 @@
# Contributor: Stephen Zhang <zsrkmyn at gmail dot com>
pkgbase=python-pytorch
-pkgname=("python-pytorch" "python-pytorch-opt" "python-pytorch-cuda" "python-pytorch-opt-cuda")
+pkgname=("python-pytorch" "python-pytorch-cuda")
_pkgname="pytorch"
-pkgver=1.9.1
-_pkgver=1.9.1
-pkgrel=2
-pkgdesc="Tensors and Dynamic neural networks in Python with strong GPU acceleration"
+pkgver=1.10.0
+_pkgver=1.10.0
+pkgrel=1
+_pkgdesc='Tensors and Dynamic neural networks in Python with strong GPU acceleration'
+pkgdesc="${_pkgdesc}"
arch=('x86_64')
url="https://pytorch.org"
license=('BSD')
@@ -40,12 +41,15 @@
"${pkgname}-kineto::git+https://github.com/pytorch/kineto"
"${pkgname}-sleef::git+https://github.com/shibatch/sleef"
"${pkgname}-onnx-tensorrt::git+https://github.com/onnx/onnx-tensorrt"
+ "${pkgname}-pocketfft::git+https://github.com/mreineck/pocketfft"
+ "${pkgname}-cudnn-frontend::git+https://github.com/NVIDIA/cudnn-frontend.git"
"${pkgname}-benchmark::git+https://github.com/google/benchmark.git"
"${pkgname}-tbb::git+https://github.com/01org/tbb"
- "${pkgname}-XNNPACK::git+https://github.com/malfet/XNNPACK.git"
+ "${pkgname}-XNNPACK::git+https://github.com/google/XNNPACK.git"
"${pkgname}-fbjni::git+https://github.com/facebookincubator/fbjni.git"
"${pkgname}-tensorpipe::git+https://github.com/pytorch/tensorpipe.git"
"${pkgname}-pybind11::git+https://github.com/pybind/pybind11.git"
+ "${pkgname}-breakpad::git+https://github.com/driazati/breakpad.git"
"${pkgname}-fbgemm::git+https://github.com/pytorch/fbgemm"
"${pkgname}-googletest::git+https://github.com/google/googletest.git"
"${pkgname}-zstd::git+https://github.com/facebook/zstd.git"
@@ -52,14 +56,14 @@
"${pkgname}-onnx::git+https://github.com/onnx/onnx.git"
"${pkgname}-protobuf::git+https://github.com/protocolbuffers/protobuf.git"
"${pkgname}-fmt::git+https://github.com/fmtlib/fmt.git"
+ https://github.com/oneapi-src/oneDNN/commit/1fe0f2594a1bfc6386fd8f6537f971d5ae9c1214.patch
+ fix_old_nnapi_lite_interpreter_config.patch
+ fix-jit-frontend-nullptr-deref.patch
fix_include_system.patch
use-system-libuv.patch
fix-building-for-torchvision.patch
- benchmark-gcc11.patch
- xnnpack-gcc11.patch
- https://github.com/pytorch/pytorch/commit/c74c0c571880df886474be297c556562e95c00e0.patch
fix_c10.patch
- disable_non_x86_64.patch)
+ 66219.patch)
sha256sums=('SKIP'
'SKIP'
'SKIP'
@@ -95,17 +99,24 @@
'SKIP'
'SKIP'
'SKIP'
+ 'SKIP'
+ 'SKIP'
+ 'SKIP'
+ '7728e99500d8034c837bbbe2b48b780d8563de4e56fff38a96766caad08cce05'
+ '21476edfa61573892a325cb8a91e13f601142e39b34e24e4575d2cdebb063b3f'
+ 'c272684a4c747f034163fcfd9dbb7264d5fe821dd25a060f0b791760ad0083ae'
'557761502bbd994d9795bef46779e4b8c60ba0b45e7d60841f477d3b7f28a00a'
'cd9ac4aaa9f946ac5eafc57cf66c5c16b3ea7ac8af32c2558fad0705411bb669'
- '689c76e89bcf403df1b4cf7ca784381967b6a6527ed6eb6d0ad6681cf789b738'
- '278fecdb45df065343f51688cc7a1665153b5189f3341a741d546b0b518eac40'
- '64833e96e47a22f88336381f25fcd73127208dc79e2074398295d88c4596c06a'
- '3d5b9d3bbba3238d8f165e582039ec07798bccc1d1f44bd91e8b1892236cb70f'
- 'ba801238afcfc58a35410e54d4ca6a638c447865c0c6b38ed16917fd6d507954'
- 'd3ef8491718ed7e814fe63e81df2f49862fffbea891d2babbcb464796a1bd680')
+ '600bd6a4bbcec9f99ab815d82cee1c2875530b2b75f4010da5ba72ce9bf31aff'
+ '4d0d7da4a3fb099ed75f3007559fad04ac96eed87c523b274fb3bb6020e6b9b8'
+ 'd86efbe915386989d75d313fc76785e6d9c5638b983f17e98cca32174ac1fcee')
+get_pyver () {
+ python -c 'import sys; print(str(sys.version_info[0]) + "." + str(sys.version_info[1]))'
+}
+
prepare() {
- cd "${_pkgname}-${pkgver}"
+ cd "${srcdir}/${_pkgname}-${pkgver}"
# generated using parse-submodules
git submodule init
@@ -124,7 +135,7 @@
git config submodule."third_party/NNPACK_deps/FP16".url "${srcdir}/${pkgname}"-FP16
git config submodule."third_party/NNPACK_deps/psimd".url "${srcdir}/${pkgname}"-psimd
git config submodule."third_party/zstd".url "${srcdir}/${pkgname}"-zstd
- git config submodule."third-party/cpuinfo".url "${srcdir}/${pkgname}"-cpuinfo
+ git config submodule."third_party/cpuinfo".url "${srcdir}/${pkgname}"-cpuinfo
git config submodule."third_party/python-enum".url "${srcdir}/${pkgname}"-enum34
git config submodule."third_party/python-peachpy".url "${srcdir}/${pkgname}"-PeachPy
git config submodule."third_party/python-six".url "${srcdir}/${pkgname}"-six
@@ -143,7 +154,10 @@
git config submodule."third_party/XNNPACK".url "${srcdir}/${pkgname}"-XNNPACK
git config submodule."third_party/fmt".url "${srcdir}/${pkgname}"-fmt
git config submodule."third_party/tensorpipe".url "${srcdir}/${pkgname}"-tensorpipe
+ git config submodule."third_party/cudnn_frontend".url "${srcdir}/${pkgname}"-cudnn-frontend
git config submodule."third_party/kineto".url "${srcdir}/${pkgname}"-kineto
+ git config submodule."third_party/pocketfft".url "${srcdir}/${pkgname}"-pocketfft
+ git config submodule."third_party/breakpad".url "${srcdir}/${pkgname}"-breakpad
git submodule update --init --recursive
@@ -156,15 +170,24 @@
# fix https://github.com/pytorch/vision/issues/3695
patch -Np1 -i "${srcdir}/fix-building-for-torchvision.patch"
- # GCC 11 fixes
- patch -Np1 -d third_party/benchmark <../benchmark-gcc11.patch
- patch -Np1 -d third_party/XNNPACK <../xnnpack-gcc11.patch
-
- # cuda 11.4 fix
- patch -Np1 <../c74c0c571880df886474be297c556562e95c00e0.patch
# cuda 11.4.1 fix
patch -Np1 -i "${srcdir}/fix_c10.patch"
+ # https://discuss.pytorch.org/t/about-build-android-sh-lite-and-nnapi/133581
+ patch -Np1 -i "${srcdir}/fix_old_nnapi_lite_interpreter_config.patch"
+
+ # fix nullptr dereference
+ patch -Np1 -i "${srcdir}/fix-jit-frontend-nullptr-deref.patch"
+
+ # disable vec tests
+ sed -e '/set(ATen_VEC_TEST_SRCS ${ATen_VEC_TEST_SRCS} PARENT_SCOPE)/d' -i aten/CMakeLists.txt
+
+ # https://github.com/pytorch/pytorch/issues/67153, https://github.com/pytorch/pytorch/pull/66219
+ patch -Np1 -i "${srcdir}/66219.patch"
+
+ # fix ideep/mkl-dnn
+ patch -Np1 -d third_party/ideep/mkl-dnn -i "${srcdir}/1fe0f2594a1bfc6386fd8f6537f971d5ae9c1214.patch"
+
# remove local nccl
rm -rf third_party/nccl/nccl
# also remove path from nccl module, so it's not checked
@@ -173,11 +196,9 @@
# fix build with google-glog 0.5 https://github.com/pytorch/pytorch/issues/58054
sed -e '/namespace glog_internal_namespace_/d' -e 's|::glog_internal_namespace_||' -i c10/util/Logging.cpp
- cd ..
+ cd "${srcdir}"
- cp -a "${_pkgname}-${pkgver}" "${_pkgname}-${pkgver}-opt"
- cp -a "${_pkgname}-${pkgver}" "${_pkgname}-${pkgver}-cuda"
- cp -a "${_pkgname}-${pkgver}" "${_pkgname}-${pkgver}-opt-cuda"
+ cp -r "${_pkgname}-${pkgver}" "${_pkgname}-${pkgver}-cuda"
export VERBOSE=1
export PYTORCH_BUILD_VERSION="${pkgver}"
@@ -203,6 +224,8 @@
export CUDAHOSTCXX=/usr/bin/g++
export CUDA_HOST_COMPILER="${CUDAHOSTCXX}"
export CUDA_HOME=/opt/cuda
+ # hide buildt-time CUDA devices
+ export CUDA_VISIBLE_DEVICES=""
export CUDNN_LIB_DIR=/usr/lib
export CUDNN_INCLUDE_DIR=/usr/include
export TORCH_NVCC_FLAGS="-Xfatbin -compress-all"
@@ -211,38 +234,24 @@
}
build() {
- echo "Building without cuda and without non-x86-64 optimizations"
+ echo "Building without cuda and with non-x86-64 optimizations"
export USE_CUDA=0
export USE_CUDNN=0
cd "${srcdir}/${_pkgname}-${pkgver}"
- echo "add_definitions(-march=x86-64)" >> cmake/MiscCheck.cmake
- patch -Np1 -i "${srcdir}/disable_non_x86_64.patch"
- python setup.py build
-
-
- echo "Building without cuda and with non-x86-64 optimizations"
- export USE_CUDA=0
- export USE_CUDNN=0
- cd "${srcdir}/${_pkgname}-${pkgver}-opt"
echo "add_definitions(-march=haswell)" >> cmake/MiscCheck.cmake
- python setup.py build
+ # this horrible hack is necessary because the current release
+ # ships inconsistent CMake which tries to build objects before
+ # thier dependencies, build twice when dependencies are available
+ python setup.py build || python setup.py build
- echo "Building with cuda and without non-x86-64 optimizations"
+ echo "Building with cuda and with non-x86-64 optimizations"
export USE_CUDA=1
export USE_CUDNN=1
cd "${srcdir}/${_pkgname}-${pkgver}-cuda"
- patch -Np1 -i "${srcdir}/disable_non_x86_64.patch"
- echo "add_definitions(-march=x86-64)" >> cmake/MiscCheck.cmake
- python setup.py build
-
-
- echo "Building with cuda and with non-x86-64 optimizations"
- export USE_CUDA=1
- export USE_CUDNN=1
- cd "${srcdir}/${_pkgname}-${pkgver}-opt-cuda"
echo "add_definitions(-march=haswell)" >> cmake/MiscCheck.cmake
- python setup.py build
+ # same horrible hack as above
+ python setup.py build || python setup.py build
}
_package() {
@@ -253,7 +262,7 @@
install -Dm644 LICENSE "${pkgdir}/usr/share/licenses/${pkgname}/LICENSE"
- pytorchpath="usr/lib/python3.9/site-packages/torch"
+ pytorchpath="usr/lib/python$(get_pyver)/site-packages/torch"
install -d "${pkgdir}/usr/lib"
# put CMake files in correct place
@@ -277,22 +286,16 @@
}
package_python-pytorch() {
+ pkgdesc+="${_pkgdesc} (with AVX2 CPU optimizations)"
+ replaces=(python-pytorch-opt)
cd "${srcdir}/${_pkgname}-${pkgver}"
_package
}
-package_python-pytorch-opt() {
- pkgdesc="Tensors and Dynamic neural networks in Python with strong GPU acceleration (with AVX2 CPU optimizations)"
- conflicts=(python-pytorch)
- provides=(python-pytorch)
-
- cd "${srcdir}/${_pkgname}-${pkgver}-opt"
- _package
-}
-
package_python-pytorch-cuda() {
- pkgdesc="Tensors and Dynamic neural networks in Python with strong GPU acceleration (with CUDA)"
+ pkgdesc="${_pkgdesc} (with CUDA and AVX2 CPU optimizations)"
depends+=(cuda cudnn magma)
+ replaces=(python-pytorch-opt-cuda)
conflicts=(python-pytorch)
provides=(python-pytorch)
@@ -300,14 +303,4 @@
_package
}
-package_python-pytorch-opt-cuda() {
- pkgdesc="Tensors and Dynamic neural networks in Python with strong GPU acceleration (with CUDA and AVX2 CPU optimizations)"
- depends+=(cuda cudnn magma)
- conflicts=(python-pytorch)
- provides=(python-pytorch python-pytorch-cuda)
-
- cd "${srcdir}/${_pkgname}-${pkgver}-opt-cuda"
- _package
-}
-
# vim:set ts=2 sw=2 et:
Deleted: benchmark-gcc11.patch
===================================================================
--- benchmark-gcc11.patch 2021-10-31 22:13:18 UTC (rev 1035134)
+++ benchmark-gcc11.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -1,30 +0,0 @@
-From 3d1c2677686718d906f28c1d4da001c42666e6d2 Mon Sep 17 00:00:00 2001
-From: Sergei Trofimovich <slyfox at inbox.ru>
-Date: Thu, 15 Oct 2020 09:12:40 +0100
-Subject: [PATCH] src/benchmark_register.h: add missing <limits> inclusion
- (#1060)
-
-Noticed missing header when was building llvm with gcc-11:
-
-```
-llvm-project/llvm/utils/benchmark/src/benchmark_register.h:17:30:
- error: 'numeric_limits' is not a member of 'std'
- 17 | static const T kmax = std::numeric_limits<T>::max();
- | ^~~~~~~~~~~~~~
-```
----
- src/benchmark_register.h | 1 +
- 1 file changed, 1 insertion(+)
-
-diff --git a/src/benchmark_register.h b/src/benchmark_register.h
-index 61377d742..204bf1d9f 100644
---- a/src/benchmark_register.h
-+++ b/src/benchmark_register.h
-@@ -1,6 +1,7 @@
- #ifndef BENCHMARK_REGISTER_H
- #define BENCHMARK_REGISTER_H
-
-+#include <limits>
- #include <vector>
-
- #include "check.h"
Deleted: disable_non_x86_64.patch
===================================================================
--- disable_non_x86_64.patch 2021-10-31 22:13:18 UTC (rev 1035134)
+++ disable_non_x86_64.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -1,15 +0,0 @@
-diff --git a/tools/setup_helpers/cmake.py b/tools/setup_helpers/cmake.py
-index d5db749d15..fd54cca6b8 100644
---- a/tools/setup_helpers/cmake.py
-+++ b/tools/setup_helpers/cmake.py
-@@ -295,6 +295,10 @@ class CMake:
- build_options.update(cmake__options)
-
- CMake.defines(args,
-+ DISABLE_AVX2=1,
-+ DISABLE_AVX512F=1,
-+ DISABLE_FMA4=1,
-+ DISABLE_SSE4=1,
- PYTHON_EXECUTABLE=sys.executable,
- PYTHON_LIBRARY=cmake_python_library,
- PYTHON_INCLUDE_DIR=distutils.sysconfig.get_python_inc(),
Modified: fix-building-for-torchvision.patch
===================================================================
--- fix-building-for-torchvision.patch 2021-10-31 22:13:18 UTC (rev 1035134)
+++ fix-building-for-torchvision.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -17,9 +17,9 @@
constexpr bool op_allowlist_check(string_view op_name) {
- assert(op_name.find("::") != string_view::npos);
+// assert(op_name.find("::") != string_view::npos);
- #if !defined(TORCH_OPERATOR_WHITELIST)
- // If the TORCH_OPERATOR_WHITELIST parameter is not defined,
- // all ops are to be registered
+ // Use assert() instead of throw() due to a gcc bug. See:
+ // https://stackoverflow.com/questions/34280729/throw-in-constexpr-function
+ // https://github.com/fmtlib/fmt/issues/682
--
2.31.1
Added: fix-jit-frontend-nullptr-deref.patch
===================================================================
--- fix-jit-frontend-nullptr-deref.patch (rev 0)
+++ fix-jit-frontend-nullptr-deref.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -0,0 +1,12 @@
+diff --color -aur pytorch-1.10.0-old/torch/csrc/jit/frontend/ir_emitter.cpp pytorch-1.10.0-new/torch/csrc/jit/frontend/ir_emitter.cpp
+--- pytorch-1.10.0-old/torch/csrc/jit/frontend/ir_emitter.cpp 2021-10-26 01:41:27.453059792 +0300
++++ pytorch-1.10.0-new/torch/csrc/jit/frontend/ir_emitter.cpp 2021-10-26 02:00:09.783068924 +0300
+@@ -1678,7 +1678,7 @@
+ << "Union type annotation `" << type_hint->repr_str()
+ << "` can hold " << vector_repr.str() << ", but none of "
+ << "those list types can hold the types of the given dict"
+- << " elements, which were unified to " << candidate->repr_str();
++ << " elements, which were unified to " << (*unified_value_type)->repr_str();
+ } else {
+ refined_type_hint = candidate;
+ }
Modified: fix_c10.patch
===================================================================
--- fix_c10.patch 2021-10-31 22:13:18 UTC (rev 1035134)
+++ fix_c10.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -7,6 +7,6 @@
static_assert(
- sizeof(void*) != sizeof(int64_t) || // if 64-bit...
+ sizeof(void*) <= sizeof(int64_t) || // if 64-bit...
- sizeof(TensorImpl) == sizeof(int64_t) * 23,
+ sizeof(TensorImpl) == sizeof(int64_t) * 24,
"You changed the size of TensorImpl on 64-bit arch."
"See Note [TensorImpl size constraints] on how to proceed.");
Added: fix_old_nnapi_lite_interpreter_config.patch
===================================================================
--- fix_old_nnapi_lite_interpreter_config.patch (rev 0)
+++ fix_old_nnapi_lite_interpreter_config.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -0,0 +1,33 @@
+# Relevant discussion: https://discuss.pytorch.org/t/about-build-android-sh-lite-and-nnapi/133581
+diff --git a/aten/src/ATen/CMakeLists.txt b/aten/src/ATen/CMakeLists.txt
+index baf9666f11..19f9a78443 100644
+--- a/aten/src/ATen/CMakeLists.txt
++++ b/aten/src/ATen/CMakeLists.txt
+@@ -130,7 +130,7 @@ add_subdirectory(quantized)
+ add_subdirectory(nnapi)
+
+ if(BUILD_LITE_INTERPRETER)
+- set(all_cpu_cpp ${generated_cpp} ${core_generated_cpp} ${cpu_kernel_cpp})
++ set(all_cpu_cpp ${generated_cpp} ${core_generated_cpp} ${ATen_NNAPI_SRCS} ${cpu_kernel_cpp})
+ append_filelist("jit_core_sources" all_cpu_cpp)
+ append_filelist("aten_cpu_source_non_codegen_list" all_cpu_cpp)
+ append_filelist("aten_native_source_non_codegen_list" all_cpu_cpp)
+diff --git a/scripts/build_android.sh b/scripts/build_android.sh
+index daad46e8fb..211f5bb429 100755
+--- a/scripts/build_android.sh
++++ b/scripts/build_android.sh
+@@ -147,7 +147,11 @@ if [ "${ANDROID_DEBUG_SYMBOLS:-}" == '1' ]; then
+ fi
+
+ if [ -n "${USE_VULKAN}" ]; then
+- CMAKE_ARGS+=("-DUSE_VULKAN=ON")
++ CMAKE_ARGS+=("-DUSE_VULKAN=${USE_VULKAN}")
++fi
++
++if [ -n "${USE_NNAPI}" ]; then
++ CMAKE_ARGS+=("-DUSE_NNAPI=${USE_NNAPI}")
+ fi
+
+ # Use-specified CMake arguments go last to allow overridding defaults
+
+
Deleted: xnnpack-gcc11.patch
===================================================================
--- xnnpack-gcc11.patch 2021-10-31 22:13:18 UTC (rev 1035134)
+++ xnnpack-gcc11.patch 2021-10-31 22:49:00 UTC (rev 1035135)
@@ -1,33 +0,0 @@
-From 042cdaf1c24c675fca5e79eb4d2665839d7df2c2 Mon Sep 17 00:00:00 2001
-From: =?UTF-8?q?Nenad=20Miks=CC=8Ca?= <nenad.miksa at microblink.com>
-Date: Mon, 3 May 2021 13:28:59 +0200
-Subject: [PATCH] GCC 11 no longer needs this polyfill
-
----
- src/xnnpack/intrinsics-polyfill.h | 6 +++---
- 1 file changed, 3 insertions(+), 3 deletions(-)
-
-diff --git a/src/xnnpack/intrinsics-polyfill.h b/src/xnnpack/intrinsics-polyfill.h
-index 3f198d88..32d7d4f3 100644
---- a/src/xnnpack/intrinsics-polyfill.h
-+++ b/src/xnnpack/intrinsics-polyfill.h
-@@ -11,8 +11,8 @@
- #if defined(__SSE2__)
- #include <emmintrin.h>
-
--// GCC any, Clang pre-8, Android NDK Clang pre-8.0.7, Apple Clang pre-11, and ICC pre-16
--#if (defined(__GNUC__) && !defined(__clang__) && !defined(__INTEL_COMPILER)) || \
-+// GCC pre-11, Clang pre-8, Android NDK Clang pre-8.0.7, Apple Clang pre-11, and ICC pre-16
-+#if (defined(__GNUC__) && !defined(__clang__) && !defined(__INTEL_COMPILER) && __GNUC__ < 11) || \
- (defined(__clang__) && !defined(__apple_build_version__) && (__clang_major__ < 8)) || \
- (defined(__clang__) && defined(__ANDROID__) && (__clang_major__ == 8) && (__clang_minor__ == 0) && (__clang_patchlevel__ < 7)) || \
- (defined(__clang__) && defined(__apple_build_version__) && (__apple_build_version__ < 11000000)) || \
-@@ -27,7 +27,7 @@ static XNN_INTRINSIC
- void _mm_storeu_si32(const void* address, __m128i v) {
- *((int*) address) = _mm_cvtsi128_si32(v);
- }
--#endif // GCC any, Clang pre-8, Android NDK Clang pre-8.0.7, Apple Clang pre-11, and ICC pre-16
-+#endif // GCC pre-11, Clang pre-8, Android NDK Clang pre-8.0.7, Apple Clang pre-11, and ICC pre-16
- #endif // SSE2
-
- #ifdef __AVX512F__
More information about the arch-commits
mailing list