diff --git a/aten/src/ATen/native/cpu/BinaryOpsKernel.cpp b/aten/src/ATen/native/cpu/BinaryOpsKernel.cpp
index 4e9c799986..bab98f6724 100644
--- a/aten/src/ATen/native/cpu/BinaryOpsKernel.cpp
+++ b/aten/src/ATen/native/cpu/BinaryOpsKernel.cpp
@@ -24,7 +24,13 @@ using namespace vec256;
 // copysign faster for the half-precision types
 template<typename T>
 T copysign(T a, T b) {
+#if (!defined(__aarch64__)) || defined(__clang__) || (__GNUC__ > 8)
+  // std::copysign gets ICE/Segfaults with gcc 7/8 on arm64
+  // (e.g. Jetson), see PyTorch PR #51834
   return std::copysign(a, b);
+#else
+  return std::signbit(b) ? -std::abs(a) : std::abs(a);
+#endif
 }
 
 // Implement copysign for half precision floats using bit ops
@@ -149,6 +155,18 @@ void div_trunc_kernel(TensorIterator& iter) {
   }
 }
 
+// this is a function because MSVC does not like us to use #if inside AT_DISPATC
+template <typename scalar_t>
+static inline scalar_t signed_zero(scalar_t sign) {
+#if (!defined(__aarch64__)) || defined(__clang__) || (__GNUC__ > 8)
+  // std::copysign gets ICE/Segfaults with gcc 7/8 on arm64
+  // (e.g. Jetson), see PyTorch PR #51834
+  return std::copysign(scalar_t(0), sign);
+#else
+  return std::signbit(sign) ? -scalar_t(0) : scalar_t(0);
+#endif
+}
+
 // NOTE: [Floor Division in Python]
 // Python's __floordiv__ operator is more complicated than just floor(a / b).
 // It aims to maintain the property: a == (a // b) * b + remainder(a, b)
@@ -201,7 +219,7 @@ void div_floor_kernel(TensorIterator& iter) {
                 floordiv += scalar_t(1.0);
               }
             } else {
-              floordiv = copysign(scalar_t(0), a / b);
+              floordiv = signed_zero(a / b);
             }
             return floordiv;
           });
diff --git a/aten/src/ATen/native/cuda/BinaryMulDivKernel.cu b/aten/src/ATen/native/cuda/BinaryMulDivKernel.cu
index e3ac2665a4..ebebcdddeb 100644
--- a/aten/src/ATen/native/cuda/BinaryMulDivKernel.cu
+++ b/aten/src/ATen/native/cuda/BinaryMulDivKernel.cu
@@ -1,10 +1,11 @@
 #include <ATen/AccumulateType.h>
 #include <ATen/Dispatch.h>
+#include <ATen/native/BinaryOps.h>
 #include <ATen/native/DispatchStub.h>
-#include <ATen/native/cuda/Loops.cuh>
 #include <ATen/native/TensorIterator.h>
-#include <ATen/native/BinaryOps.h>
 #include <c10/cuda/CUDAGuard.h>
+#include <c10/cuda/CUDAMathCompat.h>
+#include <ATen/native/cuda/Loops.cuh>
 
 // NOTE: CUDA on Windows requires that the enclosing function
 // of a __device__ lambda not have internal linkage.
@@ -139,7 +140,9 @@ void div_floor_kernel_cuda(TensorIterator& iter) {
             floordiv += scalar_t(1.0);
           }
         } else {
-          floordiv = std::copysign(scalar_t(0), a * inv_b);
+          // std::copysign gets ICE/Segfaults with gcc 7/8 on arm64
+          // (e.g. Jetson), see PyTorch PR #51834
+          floordiv = c10::cuda::compat::copysign(scalar_t(0), a * inv_b);
         }
         return floordiv;
       });
@@ -160,7 +163,9 @@ void div_floor_kernel_cuda(TensorIterator& iter) {
             floordiv += scalar_t(1.0);
           }
         } else {
-          floordiv = std::copysign(scalar_t(0), a / b);
+          // std::copysign gets ICE/Segfaults with gcc 7/8 on arm64
+          // (e.g. Jetson), see PyTorch PR #51834
+          floordiv = c10::cuda::compat::copysign(scalar_t(0), a / b);
         }
         return floordiv;
       });
diff --git a/c10/cuda/CUDAMathCompat.h b/c10/cuda/CUDAMathCompat.h
index 1fb0c3ec29..7916481ad6 100644
--- a/c10/cuda/CUDAMathCompat.h
+++ b/c10/cuda/CUDAMathCompat.h
@@ -42,11 +42,80 @@ __MATH_FUNCTIONS_DECL__ double ceil(double x) {
   return ::ceil(x);
 }
 
+__MATH_FUNCTIONS_DECL__ float fp32_from_bits(uint32_t w) {
+#if defined(__OPENCL_VERSION__)
+  return as_float(w);
+#elif defined(__CUDA_ARCH__)
+  return __uint_as_float((unsigned int)w);
+#elif defined(__INTEL_COMPILER)
+  return _castu32_f32(w);
+#else
+  union {
+    uint32_t as_bits;
+    float as_value;
+  } fp32 = {w};
+  return fp32.as_value;
+#endif
+}
+
+__MATH_FUNCTIONS_DECL__ uint32_t fp32_to_bits(float f) {
+#if defined(__OPENCL_VERSION__)
+  return as_uint(f);
+#elif defined(__CUDA_ARCH__)
+  return (uint32_t)__float_as_uint(f);
+#elif defined(__INTEL_COMPILER)
+  return _castf32_u32(f);
+#else
+  union {
+    float as_value;
+    uint32_t as_bits;
+  } fp32 = {f};
+  return fp32.as_bits;
+#endif
+}
+
+__MATH_FUNCTIONS_DECL__ double fp64_from_bits(uint64_t w) {
+#if defined(__CUDA_ARCH__)
+  return __longlong_as_double(w);
+#else
+  union {
+    uint64_t as_bits;
+    double as_value;
+  } fp64 = {w};
+  return fp64.as_value;
+#endif
+}
+
+__MATH_FUNCTIONS_DECL__ uint64_t fp64_to_bits(double f) {
+#if defined(__CUDA_ARCH__)
+  return __double_as_longlong(f);
+#else
+  union {
+    double as_value;
+    int64_t as_bits;
+  } fp64 = {f};
+  return fp64.as_bits;
+#endif
+}
+
 __MATH_FUNCTIONS_DECL__ float copysign(float x, float y) {
+#if (!defined(__aarch64__)) || defined(__clang__) || (__GNUC__ > 8)
+  // std::copysign gets ICE/Segfaults with gcc 7/8 on arm64
+  // (e.g. Jetson), see PyTorch PR #51834
   return ::copysignf(x, y);
+#else
+  return fp32_from_bits(
+      (fp32_to_bits(x) & 0x7fffffffu) | (fp32_to_bits(y) & 0x80000000u));
+#endif
 }
 __MATH_FUNCTIONS_DECL__ double copysign(double x, double y) {
+#if (!defined(__aarch64__)) || defined(__clang__) || (__GNUC__ > 8)
   return ::copysign(x, y);
+#else
+  return fp64_from_bits(
+      (fp64_to_bits(x) & 0x7fffffffffffffffull) |
+      (fp64_to_bits(y) & 0x8000000000000000ull));
+#endif
 }
 
 __MATH_FUNCTIONS_DECL__ float floor(float x) {
