Remove 9.2 related macros for CONSTEXPR (#65066)

Summary:
Removes C10_HOST_CONSTEXPR_EXCEPT_CUDA92 references in the code

Pull Request resolved: https://github.com/pytorch/pytorch/pull/65066

Reviewed By: driazati

Differential Revision: D31022520

Pulled By: janeyx99

fbshipit-source-id: f02cdc6caba5b48405575242921f5845ff18f729

diff --git a/c10/macros/Macros.h b/c10/macros/Macros.h
index 6bb3b76..8706181 100644 (file)
--- a/c10/macros/Macros.h
+++ b/c10/macros/Macros.h
@@ -394,34 +394,6 @@ __host__ __device__
 #define C10_IS_TRIVIALLY_COPYABLE(T) std::is_trivially_copyable<T>::value
 #endif
 
-// We need --expt-relaxed-constexpr in CUDA because of Eigen. This flag allows
-// device code in CUDA to call host constexpr functions. Unfortunately,
-// the CUDA compiler (at least for CUDA 9.0, 9.1 and 9.2) isn't compatible
-// with many of the constexpr things we'd like to do and the device code
-// compiler crashes when it sees one of these host-only functions.
-// It works when nvcc builds host code, but not when it builds device code
-// and notices it can call these constexpr functions from device code.
-// As a workaround, we use C10_HOST_CONSTEXPR instead of constexpr for these
-// functions. This enables constexpr when compiled on the host and applies
-// __host__ when it is compiled on the device in an attempt to stop it from
-// being called from device functions. Not sure if the latter works, but
-// even if not, it not being constexpr anymore should be enough to stop
-// it from being called from device code.
-// TODO This occurred in CUDA 9 (9.0 to 9.2). Test if this is fixed in CUDA 10.
-#if defined(__CUDA_ARCH__)
-#define C10_HOST_CONSTEXPR __host__
-#define C10_HOST_CONSTEXPR_VAR
-#else
-#define C10_HOST_CONSTEXPR constexpr
-#define C10_HOST_CONSTEXPR_VAR constexpr
-#endif
-
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ <= 9200)
-#define C10_HOST_CONSTEXPR_EXCEPT_CUDA92
-#else
-#define C10_HOST_CONSTEXPR_EXCEPT_CUDA92 constexpr
-#endif
-
 #if !defined(__clang__) && !defined(_MSC_VER) && defined(__GNUC__) && \
     __GNUC__ < 6
 #define CONSTEXPR_EXCEPT_GCC5

diff --git a/c10/util/ArrayRef.h b/c10/util/ArrayRef.h
index 21513d7..f57c67e 100644 (file)
--- a/c10/util/ArrayRef.h
+++ b/c10/util/ArrayRef.h
@@ -72,15 +72,13 @@ class ArrayRef final {
   constexpr ArrayRef(const T& OneElt) : Data(&OneElt), Length(1) {}
 
   /// Construct an ArrayRef from a pointer and length.
-  /// CUDA 9.2 fails to compile constexpr of host-only function on device
-  C10_HOST_CONSTEXPR_EXCEPT_CUDA92 ArrayRef(const T* data, size_t length)
+  C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA ArrayRef(const T* data, size_t length)
       : Data(data), Length(length) {
     debugCheckNullptrInvariant();
   }
 
   /// Construct an ArrayRef from a range.
-  /// CUDA 9.2 fails to compile constexpr of host-only function on device
-  C10_HOST_CONSTEXPR_EXCEPT_CUDA92 ArrayRef(const T* begin, const T* end)
+  C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA ArrayRef(const T* begin, const T* end)
       : Data(begin), Length(end - begin) {
     debugCheckNullptrInvariant();
   }

diff --git a/c10/util/ConstexprCrc.h b/c10/util/ConstexprCrc.h
index c896fef..5e36e46 100644 (file)
--- a/c10/util/ConstexprCrc.h
+++ b/c10/util/ConstexprCrc.h
@@ -98,7 +98,7 @@ constexpr uint64_t crc64_table[] = {
     0x29b7d047efec8728,
 };
 
-inline C10_HOST_CONSTEXPR uint64_t
+inline C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA uint64_t
 crc64impl(uint64_t accumulator, const char* data, size_t size) {
   for (size_t i = 0; i < size; ++i) {
     accumulator =
@@ -116,11 +116,12 @@ struct crc64_t final : IdWrapper<crc64_t, uint64_t> {
 };
 
 // CRC64 with Jones coefficients and an init value of 0.
-inline C10_HOST_CONSTEXPR crc64_t crc64(const char* str, size_t size) {
+inline C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA crc64_t
+crc64(const char* str, size_t size) {
   return crc64_t{detail::crc64impl(0, str, size)};
 }
 
-inline C10_HOST_CONSTEXPR crc64_t crc64(c10::string_view str) {
+inline C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA crc64_t crc64(c10::string_view str) {
   return crc64(str.data(), str.size());
 }
 } // namespace util

diff --git a/c10/util/typeid.h b/c10/util/typeid.h
index 240c69e..c651d4e 100644 (file)
--- a/c10/util/typeid.h
+++ b/c10/util/typeid.h
@@ -74,7 +74,7 @@ class C10_API TypeIdentifier final
    * is generated during run-time. Do NOT serialize the id for storage.
    */
   template <typename T>
-  static C10_HOST_CONSTEXPR TypeIdentifier Get() noexcept {
+  static C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA TypeIdentifier Get() noexcept {
     return TypeIdentifier(c10::util::get_type_index<T>());
   }
 
@@ -428,7 +428,7 @@ class C10_API TypeMeta final {
   // Below are static functions that can be called by passing a specific type.
 
   template <class T>
-  static C10_HOST_CONSTEXPR TypeIdentifier Id() noexcept {
+  static C10_HOST_CONSTEXPR_EXCEPT_WIN_CUDA TypeIdentifier Id() noexcept {
     return TypeIdentifier::Get<T>();
   }