diff --git a/aten/src/ATen/core/TensorAccessor.h b/aten/src/ATen/core/TensorAccessor.h
index 0116964e036225..95f37fcb09510c 100644
--- a/aten/src/ATen/core/TensorAccessor.h
+++ b/aten/src/ATen/core/TensorAccessor.h
@@ -1,12 +1,13 @@
 #pragma once
 
 #include <c10/macros/Macros.h>
+#include <c10/util/Deprecated.h>
 #include <stdint.h>
 #include <cstddef>
 
 namespace at {
 
-// The PtrTraits argument to the TensorAccessor/PackedTensorAccessor
+// The PtrTraits argument to the TensorAccessor/GenericPackedTensorAccessor
 // is used to enable the __restrict__ keyword/modifier for the data
 // passed to cuda.
 template <typename T>
@@ -62,7 +63,7 @@ class TensorAccessorBase {
 
 // The `TensorAccessor` is typically instantiated for CPU `Tensor`s using
 // `Tensor.accessor<T, N>()`.
-// For CUDA `Tensor`s, `PackedTensorAccessor` is used on the host and only
+// For CUDA `Tensor`s, `GenericPackedTensorAccessor` is used on the host and only
 // indexing on the device uses `TensorAccessor`s.
 template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
 class TensorAccessor : public TensorAccessorBase<T,N,PtrTraits,index_t> {
@@ -103,7 +104,7 @@ class TensorAccessor<T,1,PtrTraits,index_t> : public TensorAccessorBase<T,1,PtrT
 };
 
 
-// PackedTensorAccessorBase and PackedTensorAccessor are used on for CUDA `Tensor`s on the host
+// GenericPackedTensorAccessorBase and GenericPackedTensorAccessor are used on for CUDA `Tensor`s on the host
 // and as
 // In contrast to `TensorAccessor`s, they copy the strides and sizes on instantiation (on the host)
 // in order to transfer them on the device when calling kernels.
@@ -112,10 +113,10 @@ class TensorAccessor<T,1,PtrTraits,index_t> : public TensorAccessorBase<T,1,PtrT
 // Instantiation from data, sizes, strides is only needed on the host and std::copy isn't available
 // on the device, so those functions are host only.
 template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-class PackedTensorAccessorBase {
+class GenericPackedTensorAccessorBase {
 public:
   typedef typename PtrTraits<T>::PtrType PtrType;
-  C10_HOST PackedTensorAccessorBase(
+  C10_HOST GenericPackedTensorAccessorBase(
       PtrType data_,
       const index_t* sizes_,
       const index_t* strides_)
@@ -126,7 +127,7 @@ class PackedTensorAccessorBase {
 
   // if index_t is not int64_t, we want to have an int64_t constructor
   template <typename source_index_t, class = typename std::enable_if<std::is_same<source_index_t, int64_t>::value>::type>
-  C10_HOST PackedTensorAccessorBase(
+  C10_HOST GenericPackedTensorAccessorBase(
       PtrType data_,
       const source_index_t* sizes_,
       const source_index_t* strides_)
@@ -156,23 +157,23 @@ class PackedTensorAccessorBase {
 };
 
 template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-class PackedTensorAccessor : public PackedTensorAccessorBase<T,N,PtrTraits,index_t> {
+class GenericPackedTensorAccessor : public GenericPackedTensorAccessorBase<T,N,PtrTraits,index_t> {
 public:
   typedef typename PtrTraits<T>::PtrType PtrType;
 
-  C10_HOST PackedTensorAccessor(
+  C10_HOST GenericPackedTensorAccessor(
       PtrType data_,
       const index_t* sizes_,
       const index_t* strides_)
-      : PackedTensorAccessorBase<T, N, PtrTraits, index_t>(data_, sizes_, strides_) {}
+      : GenericPackedTensorAccessorBase<T, N, PtrTraits, index_t>(data_, sizes_, strides_) {}
 
   // if index_t is not int64_t, we want to have an int64_t constructor
   template <typename source_index_t, class = typename std::enable_if<std::is_same<source_index_t, int64_t>::value>::type>
-  C10_HOST PackedTensorAccessor(
+  C10_HOST GenericPackedTensorAccessor(
       PtrType data_,
       const source_index_t* sizes_,
       const source_index_t* strides_)
-      : PackedTensorAccessorBase<T, N, PtrTraits, index_t>(data_, sizes_, strides_) {}
+      : GenericPackedTensorAccessorBase<T, N, PtrTraits, index_t>(data_, sizes_, strides_) {}
 
   C10_DEVICE TensorAccessor<T, N - 1, PtrTraits, index_t> operator[](index_t i) {
     index_t* new_sizes = this->sizes_ + 1;
@@ -188,22 +189,22 @@ class PackedTensorAccessor : public PackedTensorAccessorBase<T,N,PtrTraits,index
 };
 
 template<typename T, template <typename U> class PtrTraits, typename index_t>
-class PackedTensorAccessor<T,1,PtrTraits,index_t> : public PackedTensorAccessorBase<T,1,PtrTraits,index_t> {
+class GenericPackedTensorAccessor<T,1,PtrTraits,index_t> : public GenericPackedTensorAccessorBase<T,1,PtrTraits,index_t> {
 public:
   typedef typename PtrTraits<T>::PtrType PtrType;
-  C10_HOST PackedTensorAccessor(
+  C10_HOST GenericPackedTensorAccessor(
       PtrType data_,
       const index_t* sizes_,
       const index_t* strides_)
-      : PackedTensorAccessorBase<T, 1, PtrTraits, index_t>(data_, sizes_, strides_) {}
+      : GenericPackedTensorAccessorBase<T, 1, PtrTraits, index_t>(data_, sizes_, strides_) {}
 
   // if index_t is not int64_t, we want to have an int64_t constructor
   template <typename source_index_t, class = typename std::enable_if<std::is_same<source_index_t, int64_t>::value>::type>
-  C10_HOST PackedTensorAccessor(
+  C10_HOST GenericPackedTensorAccessor(
       PtrType data_,
       const source_index_t* sizes_,
       const source_index_t* strides_)
-      : PackedTensorAccessorBase<T, 1, PtrTraits, index_t>(data_, sizes_, strides_) {}
+      : GenericPackedTensorAccessorBase<T, 1, PtrTraits, index_t>(data_, sizes_, strides_) {}
 
   C10_DEVICE T & operator[](index_t i) {
     return this->data_[this->strides_[0] * i];
@@ -213,4 +214,19 @@ class PackedTensorAccessor<T,1,PtrTraits,index_t> : public PackedTensorAccessorB
   }
 };
 
-}
+
+// Can't put this directly into the macro function args because of commas
+#define AT_X GenericPackedTensorAccessor<T, N, PtrTraits, index_t>
+
+// Old name for `GenericPackedTensorAccessor`
+template <typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+C10_DEFINE_DEPRECATED_USING(PackedTensorAccessor, AT_X)
+
+#undef AT_X
+
+template <typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+using PackedTensorAccessor32 = GenericPackedTensorAccessor<T, N, PtrTraits, int32_t>;
+
+template <typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+using PackedTensorAccessor64 = GenericPackedTensorAccessor<T, N, PtrTraits, int64_t>;
+} // namespace at
diff --git a/aten/src/ATen/core/TensorBody.h b/aten/src/ATen/core/TensorBody.h
index b8e45c7cc41ab0..a7c7a764f11b78 100644
--- a/aten/src/ATen/core/TensorBody.h
+++ b/aten/src/ATen/core/TensorBody.h
@@ -11,6 +11,7 @@
 #include <c10/core/TensorImpl.h>
 #include <c10/core/UndefinedTensorImpl.h>
 #include <c10/util/Exception.h>
+#include <c10/util/Deprecated.h>
 #include <c10/util/Optional.h>
 #include <c10/util/intrusive_ptr.h>
 #include <ATen/core/LegacyTypeDispatch.h>
@@ -317,19 +318,42 @@ class CAFFE2_API Tensor {
   template<typename T, size_t N>
   TensorAccessor<T,N> accessor() && = delete;
 
-  // Return a `PackedTensorAccessor` for CUDA `Tensor`s. You have to specify scalar type and
+  // Return a `GenericPackedTensorAccessor` for CUDA `Tensor`s. You have to specify scalar type and
   // dimension. You can optionally specify RestrictPtrTraits as a template parameter to
   // cast the data pointer to a __restrict__ pointer.
-  // In order to use this, your CUDA kernel has to take a corresponding PackedTensorAccessor
+  // In order to use this, your CUDA kernel has to take a corresponding GenericPackedTensorAccessor
   // as an argument.
   template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-  PackedTensorAccessor<T,N,PtrTraits,index_t> packed_accessor() const& {
+  GenericPackedTensorAccessor<T,N,PtrTraits,index_t> generic_packed_accessor() const& {
     static_assert(N > 0, "accessor is used for indexing tensor, for scalars use *data_ptr<T>()");
     TORCH_CHECK(dim() == N, "expected ", N, " dims but tensor has ", dim());
-    return PackedTensorAccessor<T,N,PtrTraits,index_t>(static_cast<typename PtrTraits<T>::PtrType>(data_ptr<T>()),sizes().data(),strides().data());
+    return GenericPackedTensorAccessor<T,N,PtrTraits,index_t>(static_cast<typename PtrTraits<T>::PtrType>(data_ptr<T>()),sizes().data(),strides().data());
   }
-  template<typename T, size_t N,  template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-  PackedTensorAccessor<T,N> packed_accessor() && = delete;
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+  GenericPackedTensorAccessor<T,N> generic_packed_accessor() && = delete;
+
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor32<T,N,PtrTraits> packed_accessor32() const& {
+    return generic_packed_accessor<T,N,PtrTraits,int32_t>();
+  }
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor32<T,N,PtrTraits> packed_accessor32() && = delete;
+
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor64<T,N,PtrTraits> packed_accessor64() const& {
+    return generic_packed_accessor<T,N,PtrTraits,int64_t>();
+  }
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor64<T,N,PtrTraits> packed_accessor64() && = delete;
+
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+  C10_DEPRECATED_MESSAGE("packed_accessor is deprecated, use packed_accessor32 or packed_accessor64 instead")
+  GenericPackedTensorAccessor<T,N,PtrTraits,index_t> packed_accessor() const & {
+    return generic_packed_accessor<T,N,PtrTraits,index_t>();
+  }
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+  C10_DEPRECATED_MESSAGE("packed_accessor is deprecated, use packed_accessor32 or packed_accessor64 instead")
+  GenericPackedTensorAccessor<T,N,PtrTraits,index_t> packed_accessor() && = delete;
 
   Tensor operator-() const;
   Tensor& operator+=(const Tensor & other);
diff --git a/aten/src/ATen/native/cuda/AveragePool3d.cu b/aten/src/ATen/native/cuda/AveragePool3d.cu
index 17fd342878871f..214e08d92bbf02 100644
--- a/aten/src/ATen/native/cuda/AveragePool3d.cu
+++ b/aten/src/ATen/native/cuda/AveragePool3d.cu
@@ -23,8 +23,8 @@ __device__ inline int max(int a, int b) {
 
 template <typename scalar_t, typename accscalar_t>
 __global__ void avg_pool3d_cuda_update_output(
-  PackedTensorAccessor<scalar_t, 4> input,
-  PackedTensorAccessor<scalar_t, 4> output,
+  PackedTensorAccessor64<scalar_t, 4> input,
+  PackedTensorAccessor64<scalar_t, 4> output,
   int kT, int kH, int kW,
   int dT, int dH, int dW,
   int padT, int padH, int padW,
@@ -87,8 +87,8 @@ __global__ void avg_pool3d_cuda_update_output(
 //
 template<int KERNEL_WIDTH, typename scalar_t, typename accscalar_t>
 __global__ void avg_pool3d_cuda_update_output(
-  PackedTensorAccessor<scalar_t, 4> input,
-  PackedTensorAccessor<scalar_t, 4> output,
+  PackedTensorAccessor64<scalar_t, 4> input,
+  PackedTensorAccessor64<scalar_t, 4> output,
   int kT, int kH,
   int dT, int dH, int dW,
   int padT, int padH, int padW,
@@ -148,8 +148,8 @@ __global__ void avg_pool3d_cuda_update_output(
 
 template <typename scalar_t, typename accscalar_t>
 __global__ void avg_pool3d_single_backward_out_frame_stride1(
-  PackedTensorAccessor<scalar_t, 4> gradOutput,
-  PackedTensorAccessor<scalar_t, 4> gradInput,
+  PackedTensorAccessor64<scalar_t, 4> gradOutput,
+  PackedTensorAccessor64<scalar_t, 4> gradInput,
   int kT, int kH, int kW,
   accscalar_t normFactor,
   int offsetZ)
@@ -193,8 +193,8 @@ __global__ void avg_pool3d_single_backward_out_frame_stride1(
 
 template <typename scalar_t, typename accscalar_t>
 __global__ void avg_pool3d_cuda_update_grad_input_atomic(
-  PackedTensorAccessor<scalar_t, 4> gradOutput,
-  PackedTensorAccessor<scalar_t, 4> gradInput,
+  PackedTensorAccessor64<scalar_t, 4> gradOutput,
+  PackedTensorAccessor64<scalar_t, 4> gradInput,
   int kT, int kH, int kW,
   int dT, int dH, int dW,
   int padT, int padH, int padW,
@@ -251,8 +251,8 @@ __global__ void avg_pool3d_cuda_update_grad_input_atomic(
 
 template <typename scalar_t, typename accscalar_t>
 __global__ void avg_pool3d_cuda_update_grad_input(
-  PackedTensorAccessor<scalar_t, 4> gradOutput,
-  PackedTensorAccessor<scalar_t, 4> gradInput,
+  PackedTensorAccessor64<scalar_t, 4> gradOutput,
+  PackedTensorAccessor64<scalar_t, 4> gradInput,
   int kT, int kH, int kW,
   int dT, int dH, int dW,
   int padT, int padH, int padW,
@@ -309,8 +309,8 @@ __global__ void avg_pool3d_cuda_update_grad_input(
 #define LAUNCH_UPDATE_OUTPUT_KERNEL_WIDTH(KW) case KW: \
   avg_pool3d_cuda_update_output<KW, scalar_t, accscalar_t>  \
     <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>( \
-       work_input.packed_accessor<scalar_t, 4>(),           \
-       work_output.packed_accessor<scalar_t, 4>(),          \
+       work_input.packed_accessor64<scalar_t, 4>(),           \
+       work_output.packed_accessor64<scalar_t, 4>(),          \
        kT, kH,                                              \
        dT, dH, dW,                                          \
        padT, padH, padW,                                    \
@@ -425,8 +425,8 @@ void avg_pool3d_out_cuda_template(
         default:
           avg_pool3d_cuda_update_output<scalar_t, accscalar_t>
             <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-               work_input.packed_accessor<scalar_t, 4>(),
-               work_output.packed_accessor<scalar_t, 4>(),
+               work_input.packed_accessor64<scalar_t, 4>(),
+               work_output.packed_accessor64<scalar_t, 4>(),
                kT, kH, kW,
                dT, dH, dW,
                padT, padH, padW,
@@ -567,8 +567,8 @@ void avg_pool3d_backward_out_cuda_template(
 
           avg_pool3d_single_backward_out_frame_stride1<scalar_t, accscalar_t>
             <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-              work_grad_output.packed_accessor<scalar_t, 4>(),
-              work_grad_input.packed_accessor<scalar_t, 4>(),
+              work_grad_output.packed_accessor64<scalar_t, 4>(),
+              work_grad_input.packed_accessor64<scalar_t, 4>(),
               kT, kH, kW,
               1.0f/divide_factor,
               offsetZ);
@@ -600,8 +600,8 @@ void avg_pool3d_backward_out_cuda_template(
           if (kernelsOverlap) {
             avg_pool3d_cuda_update_grad_input_atomic<scalar_t, accscalar_t>
               <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-                 work_grad_output.packed_accessor<scalar_t, 4>(),
-                 work_grad_input.packed_accessor<scalar_t, 4>(),
+                 work_grad_output.packed_accessor64<scalar_t, 4>(),
+                 work_grad_input.packed_accessor64<scalar_t, 4>(),
                  kT, kH, kW,
                  dT, dH, dW,
                  padT, padH, padW,
@@ -611,8 +611,8 @@ void avg_pool3d_backward_out_cuda_template(
           else {
             avg_pool3d_cuda_update_grad_input<scalar_t, accscalar_t>
               <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-                 work_grad_output.packed_accessor<scalar_t, 4>(),
-                 work_grad_input.packed_accessor<scalar_t, 4>(),
+                 work_grad_output.packed_accessor64<scalar_t, 4>(),
+                 work_grad_input.packed_accessor64<scalar_t, 4>(),
                  kT, kH, kW,
                  dT, dH, dW,
                  padT, padH, padW,
diff --git a/aten/src/ATen/native/cuda/DilatedMaxPool3d.cu b/aten/src/ATen/native/cuda/DilatedMaxPool3d.cu
index f33cced00f6997..0e9dee088897dd 100644
--- a/aten/src/ATen/native/cuda/DilatedMaxPool3d.cu
+++ b/aten/src/ATen/native/cuda/DilatedMaxPool3d.cu
@@ -20,8 +20,8 @@ __device__ inline int min(int a, int b) {
 template <typename scalar_t>
 __global__ static void max_pool3d_with_indices_single_out_frame(
   scalar_t* inputData,
-  PackedTensorAccessor<scalar_t, 4> output,
-  PackedTensorAccessor<int64_t, 4> indices,
+  PackedTensorAccessor64<scalar_t, 4> output,
+  PackedTensorAccessor64<int64_t, 4> indices,
   int itime, int iheight, int iwidth,
   int kT, int kH, int kW,
   int dT, int dH, int dW,
@@ -81,8 +81,8 @@ __global__ static void max_pool3d_with_indices_single_out_frame(
 template <int KERNEL_WIDTH, typename scalar_t>
 __global__ static void max_pool3d_with_indices_single_out_frame(
   scalar_t* inputData,
-  PackedTensorAccessor<scalar_t, 4> output,
-  PackedTensorAccessor<int64_t, 4> indices,
+  PackedTensorAccessor64<scalar_t, 4> output,
+  PackedTensorAccessor64<int64_t, 4> indices,
   int itime, int iheight, int iwidth,
   int kT, int kH,
   int dT, int dH, int dW,
@@ -143,8 +143,8 @@ __global__ static void max_pool3d_with_indices_single_out_frame(
   max_pool3d_with_indices_single_out_frame<KW>            \
   <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>( \
     input_data,                                           \
-    output.packed_accessor<scalar_t, 4>(),                \
-    indices.packed_accessor<int64_t, 4>(),                \
+    output.packed_accessor64<scalar_t, 4>(),                \
+    indices.packed_accessor64<int64_t, 4>(),                \
     itime, iheight, iwidth,                               \
     kT, kH,                                               \
     dT, dH, dW,                                           \
@@ -185,8 +185,8 @@ void max_pool3d_with_indices_out_frame(
       max_pool3d_with_indices_single_out_frame
         <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
            input_data,
-           output.packed_accessor<scalar_t, 4>(),
-           indices.packed_accessor<int64_t, 4>(),
+           output.packed_accessor64<scalar_t, 4>(),
+           indices.packed_accessor64<int64_t, 4>(),
            itime, iheight, iwidth,
            kT, kH, kW,
            dT, dH, dW,
@@ -209,8 +209,8 @@ void max_pool3d_with_indices_out_frame(
 template <typename scalar_t>
 __global__ static void max_pool3d_with_indices_backward_single_out_frame(
   scalar_t *gradInputData,
-  PackedTensorAccessor<scalar_t, 4> gradOutput,
-  PackedTensorAccessor<int64_t, 4> indices,
+  PackedTensorAccessor64<scalar_t, 4> gradOutput,
+  PackedTensorAccessor64<int64_t, 4> indices,
   int itime, int iheight, int iwidth,
   int dT, int dH, int dW,
   int pT, int pH, int pW,
@@ -255,8 +255,8 @@ void max_pool3d_with_indices_backward_out_frame(
     max_pool3d_with_indices_backward_single_out_frame
       <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
         gradInputData,
-        gradOutput.packed_accessor<scalar_t, 4>(),
-        indices.packed_accessor<int64_t, 4>(),
+        gradOutput.packed_accessor64<scalar_t, 4>(),
+        indices.packed_accessor64<int64_t, 4>(),
         itime, iheight, iwidth,
         dT, dH, dW,
         pT, pH, pW,
diff --git a/aten/src/ATen/native/cuda/FractionalMaxPool3d.cu b/aten/src/ATen/native/cuda/FractionalMaxPool3d.cu
index c44b49c004d4ee..ecd7188b273fdb 100644
--- a/aten/src/ATen/native/cuda/FractionalMaxPool3d.cu
+++ b/aten/src/ATen/native/cuda/FractionalMaxPool3d.cu
@@ -40,10 +40,10 @@ __device__ inline int64_t get_intervals(
 
 template <typename scalar_t>
 __global__ void fractional_max_pool3d_out_frame(
-  PackedTensorAccessor<scalar_t, 5> input,
-  PackedTensorAccessor<scalar_t, 5> output,
-  PackedTensorAccessor<int64_t, 5> indices,
-  PackedTensorAccessor<scalar_t, 3> samples,
+  PackedTensorAccessor64<scalar_t, 5> input,
+  PackedTensorAccessor64<scalar_t, 5> output,
+  PackedTensorAccessor64<int64_t, 5> indices,
+  PackedTensorAccessor64<scalar_t, 3> samples,
   int64_t poolSizeT, int64_t poolSizeH, int64_t poolSizeW) {
     using accscalar_t = at::acc_type<scalar_t, /*is_cuda=*/true>;
     // Output (t, h, w) point that this thread is responsible for
@@ -109,9 +109,9 @@ __global__ void fractional_max_pool3d_out_frame(
 
 template <typename scalar_t>
 __global__ void fractional_max_pool3d_backward_out_frame(
-  PackedTensorAccessor<scalar_t, 5> gradInput,
-  PackedTensorAccessor<scalar_t, 5> gradOutput,
-  PackedTensorAccessor<int64_t, 5> indices) {
+  PackedTensorAccessor64<scalar_t, 5> gradInput,
+  PackedTensorAccessor64<scalar_t, 5> gradOutput,
+  PackedTensorAccessor64<int64_t, 5> indices) {
   // Output (h, w) point that this thread is responsible for
   int64_t ourOutputPoint = threadIdx.x + blockIdx.x * blockDim.x;
   int64_t plane = blockIdx.y;
@@ -236,10 +236,10 @@ void fractional_max_pool3d_out_cuda_template(
       [&]{
         fractional_max_pool3d_out_frame<scalar_t>
         <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-          input_.packed_accessor<scalar_t, 5>(),
-          output_.packed_accessor<scalar_t, 5>(),
-          indices_.packed_accessor<int64_t, 5>(),
-          randomSamples.packed_accessor<scalar_t, 3>(),
+          input_.packed_accessor64<scalar_t, 5>(),
+          output_.packed_accessor64<scalar_t, 5>(),
+          indices_.packed_accessor64<int64_t, 5>(),
+          randomSamples.packed_accessor64<scalar_t, 3>(),
           poolSizeT, poolSizeH, poolSizeW
         );
       }
@@ -326,9 +326,9 @@ void fractional_max_pool3d_backward_out_cuda_template(
       [&] {
         fractional_max_pool3d_backward_out_frame<scalar_t>
         <<<grid, block, 0, at::cuda::getCurrentCUDAStream()>>>(
-          gradInput_.packed_accessor<scalar_t, 5>(),
-          gradOutput_.packed_accessor<scalar_t, 5>(),
-          indices_.packed_accessor<int64_t, 5>()
+          gradInput_.packed_accessor64<scalar_t, 5>(),
+          gradOutput_.packed_accessor64<scalar_t, 5>(),
+          indices_.packed_accessor64<int64_t, 5>()
         );
       }
     );
diff --git a/aten/src/ATen/native/cuda/MaxUnpooling.cu b/aten/src/ATen/native/cuda/MaxUnpooling.cu
index 1db0afd8b3afea..e4131c701bbcd5 100644
--- a/aten/src/ATen/native/cuda/MaxUnpooling.cu
+++ b/aten/src/ATen/native/cuda/MaxUnpooling.cu
@@ -38,8 +38,8 @@ __global__ void max_unpooling2d_forward_kernel(
 
 template <typename T>
 __global__ void max_unpooling3d_forward_kernel(
-    PackedTensorAccessor<T, 4> input,
-    PackedTensorAccessor<int64_t, 4> indices,
+    PackedTensorAccessor64<T, 4> input,
+    PackedTensorAccessor64<int64_t, 4> indices,
     T* output,
     const int64_t oT,
     const int64_t oH,
@@ -82,8 +82,8 @@ __global__ void max_unpooling3d_backward_kernel(
     int64_t oT,
     int64_t oH,
     int64_t oW,
-    PackedTensorAccessor<int64_t, 4> indices,
-    PackedTensorAccessor<T, 4> gradInput,
+    PackedTensorAccessor64<int64_t, 4> indices,
+    PackedTensorAccessor64<T, 4> gradInput,
     int offsetZ) {
   int iColumn = blockIdx.x * blockDim.x + threadIdx.x;
   int iRow = blockIdx.y * blockDim.y + threadIdx.y;
@@ -339,8 +339,8 @@ Tensor& max_unpooling3d_forward_out_cuda(
               block,
               0,
               at::cuda::getCurrentCUDAStream()>>>(
-              self.packed_accessor<scalar_t, 4>(),
-              indices.packed_accessor<int64_t, 4>(),
+              self.packed_accessor64<scalar_t, 4>(),
+              indices.packed_accessor64<int64_t, 4>(),
               output.data_ptr<scalar_t>(),
               oT,
               oH,
@@ -558,8 +558,8 @@ at::Tensor& max_unpooling3d_backward_out_cuda(
               oT,
               oH,
               oW,
-              indices.packed_accessor<int64_t, 4>(),
-              grad_input_reshaped.packed_accessor<scalar_t, 4>(),
+              indices.packed_accessor64<int64_t, 4>(),
+              grad_input_reshaped.packed_accessor64<scalar_t, 4>(),
               offsetZ);
           TORCH_CHECK(
               cudaGetLastError() == cudaSuccess,
diff --git a/aten/src/ATen/native/cuda/Normalization.cuh b/aten/src/ATen/native/cuda/Normalization.cuh
index cba05589d5fe88..414db1ac9e37e2 100644
--- a/aten/src/ATen/native/cuda/Normalization.cuh
+++ b/aten/src/ATen/native/cuda/Normalization.cuh
@@ -157,12 +157,12 @@ __device__ scalar_t reduce(Op op, PTA tensor, int plane) {
 
 template <typename input_scalar_t, typename stat_scalar_t, typename stat_accscalar_t, bool train, typename index_t>
 __global__ void batch_norm_transform_input_kernel(
-    const PackedTensorAccessor<input_scalar_t, 3, RestrictPtrTraits, index_t> input,
-    PackedTensorAccessor<input_scalar_t, 3, RestrictPtrTraits, index_t> output,
-    const PackedTensorAccessor<typename std::conditional<train, stat_accscalar_t, stat_scalar_t>::type, 1, RestrictPtrTraits, index_t> mean_,
-    const PackedTensorAccessor<typename std::conditional<train, stat_accscalar_t, stat_scalar_t>::type, 1, RestrictPtrTraits, index_t> var_or_invstd,
-    const PackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> weight,
-    const PackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> bias,
+    const GenericPackedTensorAccessor<input_scalar_t, 3, RestrictPtrTraits, index_t> input,
+    GenericPackedTensorAccessor<input_scalar_t, 3, RestrictPtrTraits, index_t> output,
+    const GenericPackedTensorAccessor<typename std::conditional<train, stat_accscalar_t, stat_scalar_t>::type, 1, RestrictPtrTraits, index_t> mean_,
+    const GenericPackedTensorAccessor<typename std::conditional<train, stat_accscalar_t, stat_scalar_t>::type, 1, RestrictPtrTraits, index_t> var_or_invstd,
+    const GenericPackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> weight,
+    const GenericPackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> bias,
     stat_accscalar_t epsilon) {
 
   index_t plane = blockIdx.x;
@@ -214,13 +214,13 @@ struct Var {
 
 template <template<typename T> class VarTransform, typename input_scalar_t, typename stat_scalar_t, typename stat_accscalar_t, typename index_t>
 __global__ void batch_norm_collect_statistics_kernel(
-    const PackedTensorAccessor<input_scalar_t, 3, RestrictPtrTraits, index_t> input,
+    const GenericPackedTensorAccessor<input_scalar_t, 3, RestrictPtrTraits, index_t> input,
     const stat_accscalar_t epsilon,
     const stat_accscalar_t momentum,
-    PackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> running_mean,
-    PackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> running_var,
-    PackedTensorAccessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t> save_mean,
-    PackedTensorAccessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t> save_transformed_var) {
+    GenericPackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> running_mean,
+    GenericPackedTensorAccessor<stat_scalar_t, 1, RestrictPtrTraits, index_t> running_var,
+    GenericPackedTensorAccessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t> save_mean,
+    GenericPackedTensorAccessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t> save_transformed_var) {
 
   __shared__ int shared_n[2 * 2 * C10_WARP_SIZE + C10_WARP_SIZE];
 
@@ -310,16 +310,16 @@ __global__ void batch_norm_collect_statistics_kernel(
 
 template <typename input_scalar_t, typename stat_scalar_t, typename stat_accscalar_t, typename index_t>
 __global__ void batch_norm_backward_kernel(
-    const PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> input,
-    const PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_output,
-    PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_input,
-    PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_weight,
-    PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_bias,
-    const PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> weight,
-    const PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> running_mean,
-    const PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> running_var,
-    const PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> save_mean,
-    const PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> save_invstd,
+    const GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> input,
+    const GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_output,
+    GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_input,
+    GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_weight,
+    GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_bias,
+    const GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> weight,
+    const GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> running_mean,
+    const GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> running_var,
+    const GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> save_mean,
+    const GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> save_invstd,
     bool train,
     stat_accscalar_t epsilon) {
 
@@ -341,9 +341,9 @@ __global__ void batch_norm_backward_kernel(
   // Compute two values across (batch, x/y/z) in one pass:
   // 1. Sum(grad_output)
   // 2. DotProduct(input - mean, grad_output)
-  GradOp<input_scalar_t, stat_accscalar_t, PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>> g(mean, input, grad_output);
+  GradOp<input_scalar_t, stat_accscalar_t, GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>> g(mean, input, grad_output);
   Float2<input_scalar_t, stat_accscalar_t> res = reduce<Float2<input_scalar_t, stat_accscalar_t>, GradOp<input_scalar_t, stat_accscalar_t,
-                                                                                   PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>>>(g, grad_output, plane);
+                                                                                   GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>>>(g, grad_output, plane);
   stat_accscalar_t grad_output_sum = res.v1;
   stat_accscalar_t dot_p = res.v2;
 
@@ -381,15 +381,15 @@ __global__ void batch_norm_backward_kernel(
 
 template <typename scalar_t, typename accscalar_t, typename index_t>
 __global__ void batch_norm_reduce_statistics_kernel(
-    const PackedTensorAccessor<accscalar_t, 2, RestrictPtrTraits, index_t> vec_mean,
-    const PackedTensorAccessor<accscalar_t, 2, RestrictPtrTraits, index_t> vec_invstd,
-    PackedTensorAccessor<accscalar_t, 1, RestrictPtrTraits, index_t> mean,
-    PackedTensorAccessor<accscalar_t, 1, RestrictPtrTraits, index_t> invstd,
-    PackedTensorAccessor<scalar_t, 1, RestrictPtrTraits, index_t> running_mean,
-    PackedTensorAccessor<scalar_t, 1, RestrictPtrTraits, index_t> running_var,
+    const GenericPackedTensorAccessor<accscalar_t, 2, RestrictPtrTraits, index_t> vec_mean,
+    const GenericPackedTensorAccessor<accscalar_t, 2, RestrictPtrTraits, index_t> vec_invstd,
+    GenericPackedTensorAccessor<accscalar_t, 1, RestrictPtrTraits, index_t> mean,
+    GenericPackedTensorAccessor<accscalar_t, 1, RestrictPtrTraits, index_t> invstd,
+    GenericPackedTensorAccessor<scalar_t, 1, RestrictPtrTraits, index_t> running_mean,
+    GenericPackedTensorAccessor<scalar_t, 1, RestrictPtrTraits, index_t> running_var,
     const accscalar_t epsilon,
     const accscalar_t momentum,
-    const PackedTensorAccessor<scalar_t, 1, RestrictPtrTraits, index_t> counts) {
+    const GenericPackedTensorAccessor<scalar_t, 1, RestrictPtrTraits, index_t> counts) {
 
   int feature_size = vec_mean.size(1);
   int world_size = vec_mean.size(0);
@@ -427,14 +427,14 @@ __global__ void batch_norm_reduce_statistics_kernel(
 
 template <typename input_scalar_t, typename stat_scalar_t, typename stat_accscalar_t, typename index_t>
 __global__ void batch_norm_backward_reduce_kernel(
-    const PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> input,
-    const PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_output,
-    PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean,
-    PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> invstd,
-    PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy,
-    PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy_xmu,
-    PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_weight,
-    PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_bias) {
+    const GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> input,
+    const GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_output,
+    GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean,
+    GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> invstd,
+    GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy,
+    GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy_xmu,
+    GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_weight,
+    GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> grad_bias) {
 
   index_t plane = blockIdx.x;
   index_t N = input.size(0) * input.size(2);
@@ -442,9 +442,9 @@ __global__ void batch_norm_backward_reduce_kernel(
   stat_accscalar_t r_mean = mean[plane];
   stat_accscalar_t factor = invstd[plane];
 
-  GradOp<input_scalar_t, stat_accscalar_t, PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>> g(r_mean, input, grad_output);
+  GradOp<input_scalar_t, stat_accscalar_t, GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>> g(r_mean, input, grad_output);
   Float2<input_scalar_t, stat_accscalar_t> res = reduce<Float2<input_scalar_t, stat_accscalar_t>, GradOp<input_scalar_t, stat_accscalar_t,
-                                                                                   PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>>>(g, grad_output, plane);
+                                                                                   GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t>>>(g, grad_output, plane);
 
   stat_accscalar_t norm = stat_accscalar_t(1) / N;
   if (threadIdx.x == 0) {
@@ -465,14 +465,14 @@ __global__ void batch_norm_backward_reduce_kernel(
 
 template <typename input_scalar_t, typename stat_scalar_t, typename stat_accscalar_t, typename index_t>
 __global__ void batch_norm_backward_elemt_kernel(
-    const PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> input,
-    const PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_output,
-    const PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean,
-    const PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> invstd,
-    const PackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> weight,
-    const PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy,
-    const PackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy_xmu,
-    PackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_input) {
+    const GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> input,
+    const GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_output,
+    const GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean,
+    const GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> invstd,
+    const GenericPackedTensorAccessor<stat_scalar_t, 1, DefaultPtrTraits, index_t> weight,
+    const GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy,
+    const GenericPackedTensorAccessor<stat_accscalar_t, 1, DefaultPtrTraits, index_t> mean_dy_xmu,
+    GenericPackedTensorAccessor<input_scalar_t, 3, DefaultPtrTraits, index_t> grad_input) {
 
   index_t plane = blockIdx.x;
 
@@ -502,12 +502,12 @@ __global__ void batch_norm_backward_elemt_kernel(
 }
 
 template <typename scalar_t, int64_t dim, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-static PackedTensorAccessor<scalar_t, dim, PtrTraits, index_t> packed_accessor_or_dummy(const Tensor& t) {
+static GenericPackedTensorAccessor<scalar_t, dim, PtrTraits, index_t> packed_accessor_or_dummy(const Tensor& t) {
   if (! t.defined()) {
     const std::vector<index_t> zeros(dim);
-    return PackedTensorAccessor<scalar_t, dim, PtrTraits, index_t>(nullptr, zeros.data(), zeros.data());
+    return GenericPackedTensorAccessor<scalar_t, dim, PtrTraits, index_t>(nullptr, zeros.data(), zeros.data());
   }
-  return t.packed_accessor<scalar_t, dim, PtrTraits, index_t>();
+  return t.generic_packed_accessor<scalar_t, dim, PtrTraits, index_t>();
 }
 
 template<typename input_scalar_t, typename stat_scalar_t, typename index_t>
@@ -532,7 +532,7 @@ std::tuple<Tensor, Tensor, Tensor> batch_norm_cuda_template(const Tensor& input_
 
   auto bs = input_reshaped.size(0);
   auto features = input_reshaped.size(2);
-  auto input = input_reshaped.packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
   auto input_options = input_.options();
   if (input_.scalar_type() == at::ScalarType::Half) {
     input_options = input_options.dtype(ScalarType::Float);
@@ -544,13 +544,13 @@ std::tuple<Tensor, Tensor, Tensor> batch_norm_cuda_template(const Tensor& input_
     save_mean_ = at::empty({0}, input_options);
     save_invstd_ = at::empty({0}, input_options);
   }
-  auto output = output_reshaped.packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
+  auto output = output_reshaped.generic_packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
   auto weight = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(weight_);
   auto bias = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(bias_);
   auto running_mean = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(running_mean_);
   auto running_var = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(running_var_);
-  auto save_mean = save_mean_.packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
-  auto save_invstd = save_invstd_.packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
+  auto save_mean = save_mean_.generic_packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
+  auto save_invstd = save_invstd_.generic_packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
   auto stream = at::cuda::getCurrentCUDAStream();
 
   // The input_transform kernel is pointwise, but we need to balance reading parameters (save_var/mean,
@@ -606,8 +606,8 @@ std::tuple<Tensor, Tensor, Tensor> batch_norm_backward_cuda_template(const Tenso
     grad_bias_ = at::empty_like(weight_);
   }
 
-  auto input = input_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
-  auto grad_output = grad_output_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto grad_output = grad_output_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
   auto grad_input = packed_accessor_or_dummy<input_scalar_t, 3, DefaultPtrTraits, index_t>(grad_input_reshaped);
   auto weight = packed_accessor_or_dummy<stat_scalar_t, 1, DefaultPtrTraits, index_t>(weight_);
   auto grad_weight = packed_accessor_or_dummy<stat_scalar_t, 1, DefaultPtrTraits, index_t>(grad_weight_);
@@ -643,7 +643,7 @@ std::tuple<Tensor, Tensor> batch_norm_stats_cuda_template(const Tensor& input_,
 
   auto bs = input_reshaped.size(0);
   auto features = input_reshaped.size(2);
-  auto input = input_reshaped.packed_accessor<scalar_t, 3, RestrictPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<scalar_t, 3, RestrictPtrTraits, index_t>();
   auto input_options = input_.options();
   dummy_mean_ = at::empty({0}, input_options);
   dummy_var_ = at::empty({0}, input_options);
@@ -655,8 +655,8 @@ std::tuple<Tensor, Tensor> batch_norm_stats_cuda_template(const Tensor& input_,
   invstd_ = at::empty({n_input}, input_options);
   auto mean = packed_accessor_or_dummy<accscalar_t, 1, RestrictPtrTraits, index_t>(mean_);
   auto invstd = packed_accessor_or_dummy<accscalar_t, 1, RestrictPtrTraits, index_t>(invstd_);
-  auto dummy_mean = dummy_mean_.packed_accessor<scalar_t, 1, RestrictPtrTraits, index_t>();
-  auto dummy_invstd = dummy_var_.packed_accessor<scalar_t, 1, RestrictPtrTraits, index_t>();
+  auto dummy_mean = dummy_mean_.generic_packed_accessor<scalar_t, 1, RestrictPtrTraits, index_t>();
+  auto dummy_invstd = dummy_var_.generic_packed_accessor<scalar_t, 1, RestrictPtrTraits, index_t>();
   auto stream = at::cuda::getCurrentCUDAStream();
 
   dim3 blocks(input.size(1));
@@ -680,12 +680,12 @@ Tensor batch_norm_elemt_cuda_template(const Tensor& input_, const Tensor& weight
 
   auto bs = input_reshaped.size(0);
   auto features = input_reshaped.size(2);
-  auto input = input_reshaped.packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
   auto input_options = input_.options();
   if (input_.scalar_type() == at::ScalarType::Half) {
     input_options = input_options.dtype(ScalarType::Float);
   }
-  auto output = output_reshaped.packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
+  auto output = output_reshaped.generic_packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
   auto weight = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(weight_);
   auto bias = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(bias_);
   auto mean = packed_accessor_or_dummy<stat_accscalar_t, 1, RestrictPtrTraits, index_t>(mean_);
@@ -730,8 +730,8 @@ std::tuple<Tensor, Tensor> batch_norm_gather_stats_cuda_template(const Tensor& m
   auto running_var = packed_accessor_or_dummy<scalar_t, 1, RestrictPtrTraits, index_t>(running_var_);
   auto counts = packed_accessor_or_dummy<scalar_t, 1, RestrictPtrTraits, index_t>(counts_);
 
-  auto save_mean = save_mean_.packed_accessor<accscalar_t, 1, RestrictPtrTraits, index_t>();
-  auto save_invstd = save_invstd_.packed_accessor<accscalar_t, 1, RestrictPtrTraits, index_t>();
+  auto save_mean = save_mean_.generic_packed_accessor<accscalar_t, 1, RestrictPtrTraits, index_t>();
+  auto save_invstd = save_invstd_.generic_packed_accessor<accscalar_t, 1, RestrictPtrTraits, index_t>();
   auto stream = at::cuda::getCurrentCUDAStream();
 
   int block = getNumThreads(features);
@@ -767,8 +767,8 @@ std::tuple<Tensor, Tensor, Tensor, Tensor> batch_norm_backward_reduce_cuda_templ
     grad_bias_ = at::empty({n_input}, weight_.options());
   }
 
-  auto input = input_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
-  auto grad_output = grad_output_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto grad_output = grad_output_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
   auto grad_weight = packed_accessor_or_dummy<stat_scalar_t, 1, DefaultPtrTraits, index_t>(grad_weight_);
   auto grad_bias = packed_accessor_or_dummy<stat_scalar_t, 1, DefaultPtrTraits, index_t>(grad_bias_);
   auto mean = packed_accessor_or_dummy<stat_accscalar_t, 1, DefaultPtrTraits, index_t>(mean_);
@@ -806,9 +806,9 @@ Tensor batch_norm_backward_elemt_cuda_template(const Tensor& grad_out_, const Te
   auto bs = input_reshaped.size(0);
   auto features = input_reshaped.size(2);
 
-  auto input = input_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
-  auto grad_input = grad_input_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
-  auto grad_output = grad_output_reshaped.packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto grad_input = grad_input_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
+  auto grad_output = grad_output_reshaped.generic_packed_accessor<input_scalar_t, 3, DefaultPtrTraits, index_t>();
   auto mean = packed_accessor_or_dummy<stat_accscalar_t, 1, DefaultPtrTraits, index_t>(mean_);
   auto invstd = packed_accessor_or_dummy<stat_accscalar_t, 1, DefaultPtrTraits, index_t>(invstd_);
   auto weight = packed_accessor_or_dummy<stat_scalar_t, 1, DefaultPtrTraits, index_t>(weight_);
@@ -848,11 +848,11 @@ std::tuple<Tensor, Tensor> batch_norm_update_stats_cuda_template(
   Tensor save_mean_ = at::empty({n_channels}, input_options);
   Tensor save_var_ = at::empty({n_channels}, input_options);
 
-  auto input = input_reshaped.packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
+  auto input = input_reshaped.generic_packed_accessor<input_scalar_t, 3, RestrictPtrTraits, index_t>();
   auto running_mean = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(running_mean_);
   auto running_var = packed_accessor_or_dummy<stat_scalar_t, 1, RestrictPtrTraits, index_t>(running_var_);
-  auto save_mean = save_mean_.packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
-  auto save_var = save_var_.packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
+  auto save_mean = save_mean_.generic_packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
+  auto save_var = save_var_.generic_packed_accessor<stat_accscalar_t, 1, RestrictPtrTraits, index_t>();
   auto stream = at::cuda::getCurrentCUDAStream();
 
   // for the reduction, we cannot use blocks for the batch dim, but if we have few threads in
diff --git a/aten/src/ATen/native/cuda/ReplicationPadding.cu b/aten/src/ATen/native/cuda/ReplicationPadding.cu
index c9da8f440b7297..ba51fc2105350c 100644
--- a/aten/src/ATen/native/cuda/ReplicationPadding.cu
+++ b/aten/src/ATen/native/cuda/ReplicationPadding.cu
@@ -27,8 +27,8 @@ __host__ __device__ __forceinline__ int imax(int a, int b) {
 namespace {
 template <typename scalar_t>
 __global__ void replication_pad_forward_kernel1d(
-    PackedTensorAccessor<scalar_t, 3> input,
-    PackedTensorAccessor<scalar_t, 3> output,
+    PackedTensorAccessor64<scalar_t, 3> input,
+    PackedTensorAccessor64<scalar_t, 3> output,
     int padL, int padR) {
 
   int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
@@ -50,8 +50,8 @@ __global__ void replication_pad_forward_kernel1d(
 
 template <typename scalar_t>
 __global__ void replication_pad_backward_kernel(
-    PackedTensorAccessor<scalar_t, 3> gradInput,
-    PackedTensorAccessor<scalar_t, 3> gradOutput,
+    PackedTensorAccessor64<scalar_t, 3> gradInput,
+    PackedTensorAccessor64<scalar_t, 3> gradOutput,
     int padL, int padR) {
 
   int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
@@ -73,8 +73,8 @@ __global__ void replication_pad_backward_kernel(
 
 template <typename scalar_t>
 __global__ void replication_pad_forward_kernel2d(
-    PackedTensorAccessor<scalar_t, 4> input,
-    PackedTensorAccessor<scalar_t, 4> output,
+    PackedTensorAccessor64<scalar_t, 4> input,
+    PackedTensorAccessor64<scalar_t, 4> output,
     int padT, int padB, int padL, int padR) {
 
   int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
@@ -100,8 +100,8 @@ __global__ void replication_pad_forward_kernel2d(
 
 template <typename scalar_t>
 __global__ void replication_pad_backward_kernel(
-    PackedTensorAccessor<scalar_t, 4> gradInput,
-    PackedTensorAccessor<scalar_t, 4> gradOutput,
+    PackedTensorAccessor64<scalar_t, 4> gradInput,
+    PackedTensorAccessor64<scalar_t, 4> gradOutput,
     int padT, int padB, int padL, int padR) {
 
   int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
@@ -127,8 +127,8 @@ __global__ void replication_pad_backward_kernel(
 
 template <typename scalar_t>
 __global__ void replication_pad_forward_kernel3d(
-    PackedTensorAccessor<scalar_t, 5> input,
-    PackedTensorAccessor<scalar_t, 5> output,
+    PackedTensorAccessor64<scalar_t, 5> input,
+    PackedTensorAccessor64<scalar_t, 5> output,
     int pfront, int pback, int ptop, int pbottom, int pleft, int pright) {
 
   int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
@@ -163,8 +163,8 @@ __global__ void replication_pad_forward_kernel3d(
 
 template <typename scalar_t>
 __global__ void replication_pad_backward_kernel(
-    PackedTensorAccessor<scalar_t, 5> gradInput,
-    PackedTensorAccessor<scalar_t, 5> gradOutput,
+    PackedTensorAccessor64<scalar_t, 5> gradInput,
+    PackedTensorAccessor64<scalar_t, 5> gradOutput,
     int pfront, int pback, int ptop, int pbottom, int pleft, int pright) {
   int outputPointId = threadIdx.x + blockIdx.x * blockDim.x;
   int plane = blockIdx.y;
@@ -242,8 +242,8 @@ void replication_pad1d_out_cuda_template(
         output.resize_({numPlanes, outputW});
         auto input_ = input.unsqueeze(0);
         auto output_ = output.unsqueeze(0);
-        auto devInput = input_.packed_accessor<scalar_t, 3>();
-        auto devOutput = output_.packed_accessor<scalar_t, 3>();
+        auto devInput = input_.packed_accessor64<scalar_t, 3>();
+        auto devOutput = output_.packed_accessor64<scalar_t, 3>();
 
         int outputPlaneSize = devOutput.size(2);
         dim3 gridSize(THCCeilDiv(outputPlaneSize, 256),
@@ -255,8 +255,8 @@ void replication_pad1d_out_cuda_template(
           at::cuda::getCurrentCUDAStream()>>>(devInput, devOutput, padL, padR);
       } else {
         output.resize_({numBatch, numPlanes, outputW});
-        auto devInput = input.packed_accessor<scalar_t, 3>();
-        auto devOutput = output.packed_accessor<scalar_t, 3>();
+        auto devInput = input.packed_accessor64<scalar_t, 3>();
+        auto devOutput = output.packed_accessor64<scalar_t, 3>();
 
         int outputPlaneSize = devOutput.size(2);
         dim3 gridSize(THCCeilDiv(outputPlaneSize, 256),
@@ -314,8 +314,8 @@ void replication_pad1d_backward_out_cuda_template(
       gradInput_ = gradInput.unsqueeze(0);
       gradOutput_ = gradOutput.unsqueeze(0);
       }
-      auto devGradInput = gradInput_.packed_accessor<scalar_t, 3>();
-      auto devGradOutput = gradOutput_.packed_accessor<scalar_t, 3>();
+      auto devGradInput = gradInput_.packed_accessor64<scalar_t, 3>();
+      auto devGradOutput = gradOutput_.packed_accessor64<scalar_t, 3>();
 
       int outputPlaneSize = devGradOutput.size(2);
       dim3 gridSize(THCCeilDiv(outputPlaneSize, 256),
@@ -379,8 +379,8 @@ void replication_pad2d_out_cuda_template(
         output.resize_({numPlanes, outputH, outputW});
         auto input_ = input.unsqueeze(0);
         auto output_ = output.unsqueeze(0);
-        auto devInput = input_.packed_accessor<scalar_t, 4>();
-        auto devOutput = output_.packed_accessor<scalar_t, 4>();
+        auto devInput = input_.packed_accessor64<scalar_t, 4>();
+        auto devOutput = output_.packed_accessor64<scalar_t, 4>();
 
         int outputPlaneSize = devOutput.size(2) * devOutput.size(3);
         dim3 gridSize(THCCeilDiv(outputPlaneSize, 256),
@@ -393,8 +393,8 @@ void replication_pad2d_out_cuda_template(
             devInput, devOutput, padT, padB, padL, padR);
       } else {
         output.resize_({numBatch, numPlanes, outputH, outputW});
-        auto devInput = input.packed_accessor<scalar_t, 4>();
-        auto devOutput = output.packed_accessor<scalar_t, 4>();
+        auto devInput = input.packed_accessor64<scalar_t, 4>();
+        auto devOutput = output.packed_accessor64<scalar_t, 4>();
 
         int outputPlaneSize = devOutput.size(2) * devOutput.size(3);
         dim3 gridSize(THCCeilDiv(outputPlaneSize, 256),
@@ -462,8 +462,8 @@ void replication_pad2d_backward_out_cuda_template(
           gradInput_ = gradInput.unsqueeze(0);
           gradOutput_ = gradOutput.unsqueeze(0);
         }
-        auto devGradInput = gradInput_.packed_accessor<scalar_t, 4>();
-        auto devGradOutput = gradOutput_.packed_accessor<scalar_t, 4>();
+        auto devGradInput = gradInput_.packed_accessor64<scalar_t, 4>();
+        auto devGradOutput = gradOutput_.packed_accessor64<scalar_t, 4>();
 
         int outputPlaneSize = devGradOutput.size(2) * devGradOutput.size(3);
         dim3 gridSize(THCCeilDiv(outputPlaneSize, 256),
@@ -614,8 +614,8 @@ void replication_pad3d_out_cuda_template(
         output.resize_({numPlanes, outputD, outputH, outputW});
         auto input_ = input.unsqueeze(0);
         auto output_ = output.unsqueeze(0);
-        auto devInput = input_.packed_accessor<scalar_t, 5>();
-        auto devOutput = output_.packed_accessor<scalar_t, 5>();
+        auto devInput = input_.packed_accessor64<scalar_t, 5>();
+        auto devOutput = output_.packed_accessor64<scalar_t, 5>();
 
         int outputPlaneSize = devOutput.size(2) * devOutput.size(3) *
         devOutput.size(4);
@@ -629,8 +629,8 @@ void replication_pad3d_out_cuda_template(
             devInput, devOutput, pfront, pback, ptop, pbottom, pleft, pright);
       } else {
         output.resize_({numBatch, numPlanes, outputD, outputH, outputW});
-        auto devInput = input.packed_accessor<scalar_t, 5>();
-        auto devOutput = output.packed_accessor<scalar_t, 5>();
+        auto devInput = input.packed_accessor64<scalar_t, 5>();
+        auto devOutput = output.packed_accessor64<scalar_t, 5>();
 
         int outputPlaneSize = devOutput.size(2) * devOutput.size(3) *
           devOutput.size(4);
@@ -689,8 +689,8 @@ void replication_pad3d_backward_out_cuda_template(
         gradInput_ = gradInput.unsqueeze(0);
         gradOutput_ = gradOutput.unsqueeze(0);
       }
-      auto devGradInput = gradInput_.packed_accessor<scalar_t, 5>();
-      auto devGradOutput = gradOutput_.packed_accessor<scalar_t, 5>();
+      auto devGradInput = gradInput_.packed_accessor64<scalar_t, 5>();
+      auto devGradOutput = gradOutput_.packed_accessor64<scalar_t, 5>();
 
       int outputPlaneSize = devGradOutput.size(2) * devGradOutput.size(3) *
       devGradOutput.size(4);
diff --git a/aten/src/ATen/native/cuda/UpSample.cuh b/aten/src/ATen/native/cuda/UpSample.cuh
index 3b398e27cb6e57..0bde9149136a32 100644
--- a/aten/src/ATen/native/cuda/UpSample.cuh
+++ b/aten/src/ATen/native/cuda/UpSample.cuh
@@ -166,7 +166,7 @@ __device__ __forceinline__ static int nearest_neighbor_compute_source_index(
 /* Used by UpSampleBicubic2d.cu */
 template <typename scalar_t>
 __device__ __forceinline__ static scalar_t upsample_get_value_bounded(
-    const PackedTensorAccessor<scalar_t, 4>& data,
+    const PackedTensorAccessor64<scalar_t, 4>& data,
     int batch,
     int channel,
     int height,
@@ -181,7 +181,7 @@ __device__ __forceinline__ static scalar_t upsample_get_value_bounded(
 /* Used by UpSampleBicubic2d.cu */
 template <typename scalar_t, typename accscalar_t>
 __device__ __forceinline__ static void upsample_increment_value_bounded(
-    PackedTensorAccessor<scalar_t, 4>& data,
+    PackedTensorAccessor64<scalar_t, 4>& data,
     int batch,
     int channel,
     int height,
diff --git a/aten/src/ATen/native/cuda/UpSampleBicubic2d.cu b/aten/src/ATen/native/cuda/UpSampleBicubic2d.cu
index 443e88ec078b80..cd030051728906 100644
--- a/aten/src/ATen/native/cuda/UpSampleBicubic2d.cu
+++ b/aten/src/ATen/native/cuda/UpSampleBicubic2d.cu
@@ -18,8 +18,8 @@ __global__ void upsample_bicubic2d_out_frame(
     const accscalar_t height_scale,
     const accscalar_t width_scale,
     const bool align_corners,
-    const PackedTensorAccessor<scalar_t, 4> idata,
-    PackedTensorAccessor<scalar_t, 4> odata) {
+    const PackedTensorAccessor64<scalar_t, 4> idata,
+    PackedTensorAccessor64<scalar_t, 4> odata) {
   int index = threadIdx.x + blockIdx.x * blockDim.x;
 
   const int batchsize = idata.size(0);
@@ -93,8 +93,8 @@ __global__ void upsample_bicubic2d_backward_out_frame(
     const accscalar_t height_scale,
     const accscalar_t width_scale,
     const bool align_corners,
-    PackedTensorAccessor<scalar_t, 4> idata,
-    const PackedTensorAccessor<scalar_t, 4> odata) {
+    PackedTensorAccessor64<scalar_t, 4> idata,
+    const PackedTensorAccessor64<scalar_t, 4> odata) {
   int index = threadIdx.x + blockIdx.x * blockDim.x;
 
   const int batchsize = idata.size(0);
@@ -206,8 +206,8 @@ static void upsample_bicubic2d_out_cuda_template(
       input.scalar_type(), "upsample_bicubic2d_out_frame", [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = input.packed_accessor<scalar_t, 4>();
-        auto odata = output.packed_accessor<scalar_t, 4>();
+        auto idata = input.packed_accessor64<scalar_t, 4>();
+        auto odata = output.packed_accessor64<scalar_t, 4>();
 
         // Get scaling factors
         const accscalar_t rheight = area_pixel_compute_scale<accscalar_t>(
@@ -285,8 +285,8 @@ static void upsample_bicubic2d_backward_out_cuda_template(
       grad_output.scalar_type(), "upsample_bicubic2d_backward_out_frame", [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = grad_input.packed_accessor<scalar_t, 4>();
-        auto odata = grad_output.packed_accessor<scalar_t, 4>();
+        auto idata = grad_input.packed_accessor64<scalar_t, 4>();
+        auto odata = grad_output.packed_accessor64<scalar_t, 4>();
 
         const accscalar_t rheight = area_pixel_compute_scale<accscalar_t>(
             input_height, output_height, align_corners);
diff --git a/aten/src/ATen/native/cuda/UpSampleBilinear2d.cu b/aten/src/ATen/native/cuda/UpSampleBilinear2d.cu
index 1f3f566893cc66..d8a8ed8904fa32 100644
--- a/aten/src/ATen/native/cuda/UpSampleBilinear2d.cu
+++ b/aten/src/ATen/native/cuda/UpSampleBilinear2d.cu
@@ -197,8 +197,8 @@ static void upsample_bilinear2d_out_cuda_template(
       input.scalar_type(), "upsample_bilinear2d_out_frame", [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = input.packed_accessor<scalar_t, 4>();
-        auto odata = output.packed_accessor<scalar_t, 4>();
+        auto idata = input.packed_accessor64<scalar_t, 4>();
+        auto odata = output.packed_accessor64<scalar_t, 4>();
 
         const accscalar_t rheight = area_pixel_compute_scale<accscalar_t>(
             input_height, output_height, align_corners);
diff --git a/aten/src/ATen/native/cuda/UpSampleLinear1d.cu b/aten/src/ATen/native/cuda/UpSampleLinear1d.cu
index 0f70b57344cb6e..b4fc8d5a5afd9a 100644
--- a/aten/src/ATen/native/cuda/UpSampleLinear1d.cu
+++ b/aten/src/ATen/native/cuda/UpSampleLinear1d.cu
@@ -21,8 +21,8 @@ __global__ void upsample_linear1d_out_frame(
     const int n,
     const accscalar_t rwidth,
     const bool align_corners,
-    const PackedTensorAccessor<scalar_t, 3> idata,
-    PackedTensorAccessor<scalar_t, 3> odata) {
+    const PackedTensorAccessor64<scalar_t, 3> idata,
+    PackedTensorAccessor64<scalar_t, 3> odata) {
   int index = threadIdx.x + blockIdx.x * blockDim.x;
 
   const int batchsize = idata.size(0);
@@ -70,8 +70,8 @@ __global__ void upsample_linear1d_out_frame_backward(
     const int n,
     const accscalar_t rwidth,
     const bool align_corners,
-    PackedTensorAccessor<scalar_t, 3> idata,
-    const PackedTensorAccessor<scalar_t, 3> odata) {
+    PackedTensorAccessor64<scalar_t, 3> idata,
+    const PackedTensorAccessor64<scalar_t, 3> odata) {
   int index = threadIdx.x + blockIdx.x * blockDim.x;
 
   const int batchsize = idata.size(0);
@@ -147,8 +147,8 @@ static void upsample_linear1d_out_cuda_template(
       input.scalar_type(), "upsample_linear1d_out_frame", [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = input.packed_accessor<scalar_t, 3>();
-        auto odata = output.packed_accessor<scalar_t, 3>();
+        auto idata = input.packed_accessor64<scalar_t, 3>();
+        auto odata = output.packed_accessor64<scalar_t, 3>();
 
         const accscalar_t rwidth = area_pixel_compute_scale<accscalar_t>(
           input_width, output_width, align_corners);
@@ -207,8 +207,8 @@ static void upsample_linear1d_backward_out_cuda_template(
       grad_output.scalar_type(), "upsample_linear1d_out_frame_backward", [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = grad_input.packed_accessor<scalar_t, 3>();
-        auto odata = grad_output.packed_accessor<scalar_t, 3>();
+        auto idata = grad_input.packed_accessor64<scalar_t, 3>();
+        auto odata = grad_output.packed_accessor64<scalar_t, 3>();
 
         const accscalar_t rwidth = area_pixel_compute_scale<accscalar_t>(
             input_width, output_width, align_corners);
diff --git a/aten/src/ATen/native/cuda/UpSampleTrilinear3d.cu b/aten/src/ATen/native/cuda/UpSampleTrilinear3d.cu
index 683860e8a466b7..73799b088a64e2 100644
--- a/aten/src/ATen/native/cuda/UpSampleTrilinear3d.cu
+++ b/aten/src/ATen/native/cuda/UpSampleTrilinear3d.cu
@@ -21,8 +21,8 @@ __global__ void upsample_trilinear3d_out_frame(
     const accscalar_t rheight,
     const accscalar_t rwidth,
     const bool align_corners,
-    const PackedTensorAccessor<scalar_t, 5> idata,
-    PackedTensorAccessor<scalar_t, 5> odata) {
+    const PackedTensorAccessor64<scalar_t, 5> idata,
+    PackedTensorAccessor64<scalar_t, 5> odata) {
   int index = threadIdx.x + blockIdx.x * blockDim.x;
 
   const int batchsize = idata.size(0);
@@ -105,8 +105,8 @@ __global__ void upsample_trilinear3d_backward_out_frame(
     const accscalar_t rheight,
     const accscalar_t rwidth,
     const bool align_corners,
-    PackedTensorAccessor<scalar_t, 5> idata,
-    const PackedTensorAccessor<scalar_t, 5> odata) {
+    PackedTensorAccessor64<scalar_t, 5> idata,
+    const PackedTensorAccessor64<scalar_t, 5> odata) {
   int index = threadIdx.x + blockIdx.x * blockDim.x;
 
   const int batchsize = idata.size(0);
@@ -245,8 +245,8 @@ static void upsample_trilinear3d_out_cuda_template(
       input.scalar_type(), "upsample_trilinear3d_out_frame", [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = input.packed_accessor<scalar_t, 5>();
-        auto odata = output.packed_accessor<scalar_t, 5>();
+        auto idata = input.packed_accessor64<scalar_t, 5>();
+        auto odata = output.packed_accessor64<scalar_t, 5>();
 
         const accscalar_t rdepth = area_pixel_compute_scale<accscalar_t>(
             input_depth, output_depth, align_corners);
@@ -332,8 +332,8 @@ static void upsample_trilinear3d_backward_out_cuda_template(
       [&] {
         using accscalar_t = at::acc_type<scalar_t, true>;
 
-        auto idata = grad_input.packed_accessor<scalar_t, 5>();
-        auto odata = grad_output.packed_accessor<scalar_t, 5>();
+        auto idata = grad_input.packed_accessor64<scalar_t, 5>();
+        auto odata = grad_output.packed_accessor64<scalar_t, 5>();
 
         const accscalar_t rdepth = area_pixel_compute_scale<accscalar_t>(
             input_depth, output_depth, align_corners);
diff --git a/aten/src/ATen/templates/TensorBody.h b/aten/src/ATen/templates/TensorBody.h
index 49c23e695fda72..1fd7c4e16542f6 100644
--- a/aten/src/ATen/templates/TensorBody.h
+++ b/aten/src/ATen/templates/TensorBody.h
@@ -11,6 +11,7 @@
 #include <c10/core/TensorImpl.h>
 #include <c10/core/UndefinedTensorImpl.h>
 #include <c10/util/Exception.h>
+#include <c10/util/Deprecated.h>
 #include <c10/util/Optional.h>
 #include <c10/util/intrusive_ptr.h>
 #include <ATen/core/LegacyTypeDispatch.h>
@@ -317,19 +318,42 @@ class CAFFE2_API Tensor {
   template<typename T, size_t N>
   TensorAccessor<T,N> accessor() && = delete;
 
-  // Return a `PackedTensorAccessor` for CUDA `Tensor`s. You have to specify scalar type and
+  // Return a `GenericPackedTensorAccessor` for CUDA `Tensor`s. You have to specify scalar type and
   // dimension. You can optionally specify RestrictPtrTraits as a template parameter to
   // cast the data pointer to a __restrict__ pointer.
-  // In order to use this, your CUDA kernel has to take a corresponding PackedTensorAccessor
+  // In order to use this, your CUDA kernel has to take a corresponding GenericPackedTensorAccessor
   // as an argument.
   template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-  PackedTensorAccessor<T,N,PtrTraits,index_t> packed_accessor() const& {
+  GenericPackedTensorAccessor<T,N,PtrTraits,index_t> generic_packed_accessor() const& {
     static_assert(N > 0, "accessor is used for indexing tensor, for scalars use *data_ptr<T>()");
     TORCH_CHECK(dim() == N, "expected ", N, " dims but tensor has ", dim());
-    return PackedTensorAccessor<T,N,PtrTraits,index_t>(static_cast<typename PtrTraits<T>::PtrType>(data_ptr<T>()),sizes().data(),strides().data());
+    return GenericPackedTensorAccessor<T,N,PtrTraits,index_t>(static_cast<typename PtrTraits<T>::PtrType>(data_ptr<T>()),sizes().data(),strides().data());
   }
-  template<typename T, size_t N,  template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
-  PackedTensorAccessor<T,N> packed_accessor() && = delete;
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+  GenericPackedTensorAccessor<T,N> generic_packed_accessor() && = delete;
+
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor32<T,N,PtrTraits> packed_accessor32() const& {
+    return generic_packed_accessor<T,N,PtrTraits,int32_t>();
+  }
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor32<T,N,PtrTraits> packed_accessor32() && = delete;
+
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor64<T,N,PtrTraits> packed_accessor64() const& {
+    return generic_packed_accessor<T,N,PtrTraits,int64_t>();
+  }
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits>
+  PackedTensorAccessor64<T,N,PtrTraits> packed_accessor64() && = delete;
+
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+  C10_DEPRECATED_MESSAGE("packed_accessor is deprecated, use packed_accessor32 or packed_accessor64 instead")
+  GenericPackedTensorAccessor<T,N,PtrTraits,index_t> packed_accessor() const & {
+    return generic_packed_accessor<T,N,PtrTraits,index_t>();
+  }
+  template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
+  C10_DEPRECATED_MESSAGE("packed_accessor is deprecated, use packed_accessor32 or packed_accessor64 instead")
+  GenericPackedTensorAccessor<T,N,PtrTraits,index_t> packed_accessor() && = delete;
 
   Tensor operator-() const;
   Tensor& operator+=(const Tensor & other);
diff --git a/aten/src/ATen/test/cuda_packedtensoraccessor_test.cu b/aten/src/ATen/test/cuda_packedtensoraccessor_test.cu
index ff0c0c4eb89e5e..12d3b3d9731f1b 100644
--- a/aten/src/ATen/test/cuda_packedtensoraccessor_test.cu
+++ b/aten/src/ATen/test/cuda_packedtensoraccessor_test.cu
@@ -9,9 +9,9 @@
 using namespace at;
 
 __global__ void test_tensor_packed_accessor_kernel(
-    PackedTensorAccessor<float, 1, RestrictPtrTraits> resa,
-    PackedTensorAccessor<float, 2, RestrictPtrTraits> t1a,
-    PackedTensorAccessor<float, 1, RestrictPtrTraits> t2a) {
+    PackedTensorAccessor64<float, 1, RestrictPtrTraits> resa,
+    PackedTensorAccessor64<float, 2, RestrictPtrTraits> t1a,
+    PackedTensorAccessor64<float, 1, RestrictPtrTraits> t2a) {
   for (int64_t i = 0; i < resa.size(0); i++) {
     float val = 0.0f;
     for (int64_t j = 0; j < t1a.size(1); j++) {
@@ -21,7 +21,7 @@ __global__ void test_tensor_packed_accessor_kernel(
   }
 }
 
-// test PackedTensorAccessor and Tensor.packed_accessor
+// test GenericPackedTensorAccessor and Tensor.generic_packed_accessor
 TEST(PackedtensoraccessorTest, PackedtensoraccessorTestCUDA) {
   if (!at::cuda::is_available()) return;
   manual_seed(123);
@@ -30,9 +30,9 @@ TEST(PackedtensoraccessorTest, PackedtensoraccessorTestCUDA) {
   Tensor t2 = rand({4}, CUDA(kFloat));
   Tensor res = empty({4}, CUDA(kFloat));
 
-  auto t1a = t1.packed_accessor<float, 2, RestrictPtrTraits>();
-  auto t2a = t2.packed_accessor<float, 1, RestrictPtrTraits>();
-  auto resa = res.packed_accessor<float, 1, RestrictPtrTraits>();
+  auto t1a = t1.packed_accessor64<float, 2, RestrictPtrTraits>();
+  auto t2a = t2.packed_accessor64<float, 1, RestrictPtrTraits>();
+  auto resa = res.packed_accessor64<float, 1, RestrictPtrTraits>();
 
   auto stream = at::cuda::getCurrentCUDAStream();
 
diff --git a/docs/cpp/source/notes/tensor_basics.rst b/docs/cpp/source/notes/tensor_basics.rst
index 5d25efcf68de95..09032546a3a9ae 100644
--- a/docs/cpp/source/notes/tensor_basics.rst
+++ b/docs/cpp/source/notes/tensor_basics.rst
@@ -76,20 +76,25 @@ CUDA accessors
 .. code-block:: cpp
 
   __global__ void packed_accessor_kernel(
-      PackedTensorAccessor<float, 2> foo,
+      PackedTensorAccessor64<float, 2> foo,
       float* trace) {
     int i=threadIdx.x
     atomicAdd(trace, foo[i][i])
   }
-  
+
   torch::Tensor foo = torch::rand({12, 12});
 
   // assert foo is 2-dimensional and holds floats.
-  auto foo_a = foo.packed_accessor<float,2>();
+  auto foo_a = foo.packed_accessor64<float,2>();
   float trace = 0;
 
   packed_accessor_kernel<<<1, 12>>>(foo_a, &trace);
 
+In addition to ``PackedTensorAccessor64`` and ``packed_accessor64`` there are
+also the corresponding ``PackedTensorAccessor32`` and ``packed_accessor32``
+which use 32-bit integers for indexing. This can be quite a bit faster on CUDA
+but may lead to overflows in the indexing calculations.
+
 Note that the template can hold other parameters such as the pointer restriction
 and the integer type for indexing. See documentation for a thorough template
 description of *accessors* and *packed accessors*.