Update

malfet · malfet · commit 40a0c8c9170a · 2025-08-03T12:23:22.000-07:00
[ghstack-poisoned]
diff --git a/c10/metal/atomic.h b/c10/metal/atomic.h
@@ -85,15 +85,13 @@ struct AtomicType<uchar> {
   }
 };
 
-#if __METAL_VERSION__ >= 310
 template <>
 struct AtomicType<bfloat> {
   using type = ::metal::atomic<uint>;
   static inline void atomic_add(device type* data, long offset, bfloat value) {
     atomic_add_helper<bfloat>(data, offset, value);
   }
 };
-#endif
 
 // Metal supports atomic_store_explicit for bools, but
 // sizeof(::metal::atomic_bool) is 4 Therefore it could not be used to
diff --git a/c10/metal/common.h b/c10/metal/common.h
@@ -9,7 +9,6 @@
 #define C10_METAL_CONSTEXPR constexpr
 #endif
 
-#if !defined(__METAL__) || __METAL_VERSION__ >= 310
 #define C10_METAL_ALL_TYPES_FUNCTOR(_) \
   _(Byte, 0)                           \
   _(Char, 1)                           \
@@ -22,19 +21,6 @@
   _(ComplexFloat, 9)                   \
   _(Bool, 11)                          \
   _(BFloat16, 15)
-#else
-#define C10_METAL_ALL_TYPES_FUNCTOR(_) \
-  _(Byte, 0)                           \
-  _(Char, 1)                           \
-  _(Short, 2)                          \
-  _(Int, 3)                            \
-  _(Long, 4)                           \
-  _(Half, 5)                           \
-  _(Float, 6)                          \
-  _(ComplexHalf, 8)                    \
-  _(ComplexFloat, 9)                   \
-  _(Bool, 11)
-#endif
 
 namespace c10 {
 namespace metal {
diff --git a/c10/metal/indexing.h b/c10/metal/indexing.h
@@ -186,10 +186,8 @@ inline T val_at_offs(constant void* ptr, long offs, ScalarType type) {
       return cast_to<T>(val_at_offs<float>(ptr, offs));
     case ScalarType::Half:
       return cast_to<T>(val_at_offs<half>(ptr, offs));
-#if __METAL_VERSION__ >= 310
     case ScalarType::BFloat16:
       return cast_to<T>(val_at_offs<bfloat>(ptr, offs));
-#endif
       // Complex
     case ScalarType::ComplexHalf:
       return cast_to<T>(val_at_offs<half2>(ptr, offs));
diff --git a/c10/metal/reduction_utils.h b/c10/metal/reduction_utils.h
@@ -15,12 +15,10 @@ struct simd_type {
 template <typename T>
 using simd_type_t = typename simd_type<T>::t;
 
-#if __METAL_VERSION__ >= 310
 template <>
 struct simd_type<bfloat> {
   using t = float;
 };
-#endif
 } // namespace detail
 
 template <typename T>
diff --git a/c10/metal/utils.h b/c10/metal/utils.h
@@ -24,14 +24,12 @@ struct vectypes<half> {
   using type2 = half2;
 };
 
-#if __METAL_VERSION__ >= 310
 template <>
 struct vectypes<bfloat> {
   using type4 = bfloat4;
   using type3 = bfloat3;
   using type2 = bfloat2;
 };
-#endif
 
 template <>
 struct vectypes<short> {
@@ -79,12 +77,10 @@ struct OpMathType<uchar> {
   using type = int;
 };
 
-#if __METAL_VERSION__ >= 310
 template <>
 struct OpMathType<bfloat> {
   using type = float;
 };
-#endif
 
 // Type promotion structure for higher precision accumulation
 template <typename T>
@@ -98,13 +94,11 @@ struct AccumulationType<half> {
   using type = float;
 };
 
-#if __METAL_VERSION__ >= 310
 // Specialization for bfloat - promote to float for accumulation
 template <>
 struct AccumulationType<bfloat> {
   using type = float;
 };
-#endif
 
 } // namespace detail
 
@@ -130,7 +124,6 @@ min(T a, U b) {
   return ::metal::min(a, static_cast<T>(b));
 }
 
-#if __METAL_VERSION__ >= 310
 template <>
 inline bfloat min(bfloat a, bfloat b) {
   return bfloat(
@@ -142,7 +135,6 @@ inline bfloat max(bfloat a, bfloat b) {
   return bfloat(
       ::metal::isunordered(a, b) ? NAN : ::metal::max(float(a), float(b)));
 }
-#endif
 
 template <typename T>
 using vec2type_t = typename detail::vectypes<T>::type2;
