jhlee525 opened a new pull request, #18623:
URL: https://github.com/apache/tvm/pull/18623
This PR adds bfloat16 support for Metal , enabling efficient bfloat16
simdgroup matrix operations using Metal 3.1+.
## Summary
* Full BF16 pipeline support for Metal simdgroup matrix operations
* Runtime capability detection with safe fallback behavior
* Correct bf16 handling across scheduling, codegen, and intrinsic lowering
* Metal language version selection aligned with target capabilities
## Changes
### `Target` update
* Introduces a `min_metal_version` attribute for the Metal target
* Adds `supports_bf16()` helper to detect bf16 availability based on
`min_metal_version`
### SIMDgroup bf16 Intrinsics
Registers BF16-specific simdgroup matrix intrinsics (8×8×8):
* `simdgroup_load_8x8x8_bf16_shared`
* `simdgroup_load_8x8x8_bf16_shared_trans`
* `simdgroup_store_8x8x8_bf16_global`
* `simdgroup_store_8x8x8_bf16_shared`
* `simdgroup_make_filled_8x8x8_bf16` (renamed for consistency)
* `simdgroup_multiply_accumulate_8x8x8_bf16`
### Scheduling Updates
Updates `MetalMatmul` (in dlight) to:
* Detect input dtypes & verify bf16 support
* Apply SIMD intrinsics
* Gracefully fall back to generic scheduling with improved error handling
### Codegen Enhancements
* Fixes BF16 type handling in simdgroup codegen:
* Cast immediate values to the correct dtype in
`make_filled_simdgroup_matrix`
* Emit explicit `bfloat()` casts for bf16 immediates (Metal does not
support an `h` suffix for bf16 literals)
* Uses Metal 3.1 language version when bf16 support is available, ensuring
correct feature enablement and compatibility
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
