-
Notifications
You must be signed in to change notification settings - Fork 548
MetalPerformancePrimitives macOS xcode26.0 b6
Rolf Bjarne Kvinge edited this page Aug 19, 2025
·
2 revisions
#MetalPerformancePrimitives.framework https://github.com/dotnet/macios/issues/23418
diff -ruN /Applications/Xcode_26.0.0-beta5.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h /Applications/Xcode_26.0.0-beta6.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h
--- /Applications/Xcode_26.0.0-beta5.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h 2025-07-26 04:24:17
+++ /Applications/Xcode_26.0.0-beta6.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h 2025-08-11 05:10:54
@@ -418,6 +418,10 @@
"Scope template argument should be of op_scope type");
public:
+
+ static constexpr constant matmul2d_descriptor descriptor = Descriptor;
+ using scope = Scope;
+
matmul2d() thread = default;
template <
diff -ruN /Applications/Xcode_26.0.0-beta5.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsConvolution2dImpl.h /Applications/Xcode_26.0.0-beta6.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsConvolution2dImpl.h
--- /Applications/Xcode_26.0.0-beta5.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsConvolution2dImpl.h 2025-07-26 07:32:07
+++ /Applications/Xcode_26.0.0-beta6.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsConvolution2dImpl.h 2025-08-11 05:10:54
@@ -4257,7 +4257,10 @@
const thread convolution2d_descriptor &,
__tensor_ops_detail::__thread_void_t, uint16_t,
__tensor_ops_detail::__tensor_ops_datatype, int threads);
-
+extern "C" TENSOROPS_EXPORT EXTERNALLY_DEFINED_ATTR uint16_t
+__tensorops_impl_conv2d_cooperative_destination_tensor_get_element_index(
+ const thread convolution2d_descriptor &, __tensor_ops_detail::__thread_void_t,
+ __tensor_ops_detail::__const_thread_void_t, __tensor_ops_detail::__tensor_ops_datatype, int threads);
extern "C" EXTERNALLY_DEFINED_ATTR void
__tensorops_impl_convolution2d_op_cooperative_destination_tensor_load_dv_f32(
thread convolution2d_descriptor &desc, thread void *storage,
@@ -4701,20 +4704,7 @@
{
metal::execution_threads t = scope();
int threads = t.size();
- __tensor_ops_detail::__tensor_ops_datatype dataType;
- if constexpr (__tensor_ops_detail::__is_same_v<element_t, int32_t>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float32;
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, float>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float32;
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, half>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float16;
-#if __HAVE_BFLOAT__
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, bfloat>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_bfloat16;
-#endif
- else
- static_assert(__tensor_ops_detail::__assert_false_v<element_type>,
- "unsupported data type");
+ __tensor_ops_detail::__tensor_ops_datatype dataType = __tensor_ops_detail::__type_to_tensor_ops_datatype<element_t>::value;
return (thread element_t *)
__tensorops_impl_conv2d_cooperative_destination_tensor_elements(
@@ -4722,28 +4712,19 @@
}
static index_t get_element_index(thread_storage_t storage,
- const thread element_type *) {
- // TODO
+ const thread element_type *element) {
+
+ metal::execution_threads t = scope();
+ int threads = t.size();
+ __tensor_ops_detail::__tensor_ops_datatype dataType = __tensor_ops_detail::__type_to_tensor_ops_datatype<element_t>::value;
+ return __tensorops_impl_conv2d_cooperative_destination_tensor_get_element_index(descriptor, (__tensor_ops_detail::__thread_void_t)storage, (const thread void*)element, dataType, threads);
}
static bool is_valid_element(const_thread_storage_t storage, index_t idx)
{
metal::execution_threads t = scope();
int threads = t.size();
- __tensor_ops_detail::__tensor_ops_datatype dataType;
- if constexpr (__tensor_ops_detail::__is_same_v<element_t, int32_t>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float32;
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, float>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float32;
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, half>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float16;
-#if __HAVE_BFLOAT__
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, bfloat>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_bfloat16;
-#endif
- else
- static_assert(__tensor_ops_detail::__assert_false_v<element_type>,
- "unsupported data type");
+ __tensor_ops_detail::__tensor_ops_datatype dataType = __tensor_ops_detail::__type_to_tensor_ops_datatype<element_t>::value;
return __tensorops_impl_conv2d_cooperative_destination_tensor_is_valid_element(
descriptor, (__tensor_ops_detail::__thread_void_t)storage, idx,
@@ -4760,21 +4741,7 @@
static_assert(rank == 4, "multidimensional_indices returns 4D indices");
- __tensor_ops_detail::__tensor_ops_datatype dataType;
- if constexpr (__tensor_ops_detail::__is_same_v<element_t, int32_t>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float32;
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, float>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float32;
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, half>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_float16;
-#if __HAVE_BFLOAT__
- else if constexpr (__tensor_ops_detail::__is_same_v<element_t, bfloat>)
- dataType = __tensor_ops_detail::__tensor_ops_datatype_bfloat16;
-#endif
- else
- static_assert(__tensor_ops_detail::__assert_false_v<element_type>,
- "unsupported data type");
-
+ __tensor_ops_detail::__tensor_ops_datatype dataType = __tensor_ops_detail::__type_to_tensor_ops_datatype<element_t>::value;
if constexpr (__tensor_ops_detail::__is_same_v<coord_t, ushort>)
{
ushort coords[4];
diff -ruN /Applications/Xcode_26.0.0-beta5.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsUtility.h /Applications/Xcode_26.0.0-beta6.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsUtility.h
--- /Applications/Xcode_26.0.0-beta5.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsUtility.h 2025-07-26 07:01:33
+++ /Applications/Xcode_26.0.0-beta6.app/Contents/Developer/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/__impl/MPPTensorOpsUtility.h 2025-08-11 05:10:54
@@ -33,6 +33,8 @@
template <typename T>
struct __type_to_tensor_ops_datatype
{
+ static_assert(__tensor_ops_detail::__assert_false_v<T>,
+ "unsupported data type");
};
template <>
struct __type_to_tensor_ops_datatype<float>