4 ms·

I have questions! What is the size of the matrix elements? My understanding is that float32 is super-easy to represent in SPIR-V, and that float16 is possible

I have questions! What is the size of the matrix elements? My understanding is that float32 is super-easy to represent in SPIR-V, and that float16 is possible but the path is not as smoothly paved. Of course, on the Metal side "half" is fine. Does the M1 silicon have anything like cooperative matrices[1]? This is a huge performance bump, but currently requires fiddly vendor-specific extensions. Hopefully a vendor-neutral standard will emerge before long. Conversely, are cooperative matrices used in the A100 comparison? Are you using MoltenVK to run the dispatches? I've had ok results with this but ultimately built my own GPU API abstraction layer, so my code is calling into Metal. There are certain features, such as access to "memoryless" buffers, that may be difficult to access through MoltenVK, as I believe the concept doesn't really exist in Vulkan yet. Is that something that might be useful in a machine learning context? Do you support operations like prefix sums? If so, are these using the decoupled look-back algorithm, or multiple dispatches? I'm excited to see more portable use of compute shaders, I think to a large extent it's the future. [1]: https://developer.nvidia.com/blog/machine-learning-acceleration-vulkan-cooperative-matrices/ https://developer.nvidia.com/blog/machine-learning-accelerat...

magic_at_enimai 5y ago
Here are the matmul sizes for the MiniLM model used for inference: https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/bert_minilm_l12_h384_matmul.txt https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/b... These are the matmul sizes for the BERT training workload https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/benchmark_small_sizes.txt https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/b... Yes we use the latest MoltenVK (1.3.204.0) installed in the system. I will let @noxa and other IREE devs chime in on the SPIR-V path but we do support prefix sums etc in the GPU path. //part of nod.ai team.
- raphlinus 5y ago
  Thanks for the matmul sizes, but the question I am more interested in is precision. Matrix multiply throughput can be dramatically higher for tensor cores[1] than normal shader ALU, specifically for reduced precision arithmetic. I'm wondering to what extent that's accessible on the M1, and to what extent IREE can address them. Regarding prefix sum, the specific question I'm interested in is that SPIRV OpControlBarrier with device scope gets translated into threadgroup_barrier(mem_device) [2]. That's insufficient to make decoupled look-back work. Conversely, if you're not using decoupled look-back, you're not getting the full throughput on GPUs that do support that barrier. I'm wondering how your infrastructure deals with that. [1]: https://developer.nvidia.com/blog/programming-tensor-cores-cuda-9/ https://developer.nvidia.com/blog/programming-tensor-cores-c... [2]: https://github.com/linebender/piet-gpu/blob/d81e5cb4ee145abd1f2f01904579cb3f93960e97/tests/shader/gen/prefix.msl#L136 https://github.com/linebender/piet-gpu/blob/d81e5cb4ee145abd...
  - magic_at_enimai 5y ago
    So with Tensorcores you use TF32 which is more like FP19-ish and the marketing makes you think you get 8x the performance. But if you want actual FP32 precision you will need something like [1] but then your performance in the Tensorcore path is _only_ 2X faster than the SIMT path. I'll leave the prefix sum for other devs who know more :D https://github.com/NVIDIA/cutlass/blob/master/examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm/27_ampere_3xtf32_fast_accurate_tensorop_gemm.cu https://github.com/NVIDIA/cutlass/blob/master/examples/27_am... //part of nod.ai/shark team
    - raphlinus 5y ago
      I think we're talking past each other to some extent. Putting aside the question of how misleading it is to market a 16 bit multiply as a "TF32" operation, this is all about tradeoffs. The specific tradeoff that these tensor cores make is that in exchange for reduced precision (and a programming model which is even more of a pain than ordinary compute shaders, an astonishing achievement in and of itself), you get a lot more throughput. For certain AI workloads, particularly inference, that tradeoff is well worth it. Reading between the lines a little, it sounds like your infrastructure is potentially able to exploit a good deal of the available throughput for FP32 workloads. That's great, and I'm happy to see it! However, for workloads that don't need that much precision, the tradeoff might be a lot less advantageous to M1. That may change again if and when Apple opens up lower-level APIs to their hardware, or reverse engineering delivers usable results.
      - shaklee3 5y ago
        tf32 and fp16 tensor cores are completely different, and tf32 is not 16 bit multiplication.
noxa 5y ago
RE cooperative matrix: they have functional units (AMX/ANE) that could hopefully be exposed in MSL via something shaped like cooperative matrix and I'm pretty sure it'd be fantastic. Everything today is locked behind CoreML and Accelerate and those are poor targets for modern compiler-based approaches :( On the Vulkan side there's been rumblings of a vendor-agnostic extension for cooperative matrix and support from major vendors - at which point I'm hoping that leads to Apple wanting to show off their own HW features. RE MoltenVK: we were surprised how robust it is nowadays - we're definitely going to build out our own Metal backend for our abstraction layer but wanted to see what we could hit with the zero-code option and it's proven very useful for that! A MoltenVK build is on the order of ~12MB (last I looked) while the entire IREE runtime is 50-100KB so it's a hard pill to swallow, all other issues (security/memory consumption/startup time/etc) notwithstanding :) AFAIK the memoryless storage is only for textures and mostly useful in render passes - Vulkan has this via VK_MEMORY_PROPERTY_LAZILY_ALLOCATED_BIT. In a way what we do in our compiler is fuse dispatches such that they never hit memory at all so it hasn't been something we've needed yet. For heavy compute workloads more control over virtual memory would be useful, though, ala https://developer.nvidia.com/blog/introducing-low-level-gpu-virtual-memory-management/ https://developer.nvidia.com/blog/introducing-low-level-gpu-... - some of the biggest performance hurdles when it comes to highly dynamic ML is around memory management (being able to zero-copy resize buffers when data-dependent shapes change would be killer). // IREE dev
- raphlinus 5y ago
  Great, thanks. That answers my questions. I'll read up on the lazily allocated bit; I wasn't aware that this provided similar functionality as dispatchThreadsPerTile[1], but perhaps it's something I'm misunderstanding. I'm excited about that as a way to stitch 2D graphics rendering operations together without having to hit main memory, but from your explanation I can see that functionality might not be very useful in AI workloads. Amen on more control over dynamic memory access patterns. It's something I'm struggling with too, and I have a feeling that whatever solution I come up with is going to be a compromise. Keep up the good work, these are exciting times! [1]: https://developer.apple.com/documentation/metal/mtlrendercommandencoder/2866171-dispatchthreadspertile https://developer.apple.com/documentation/metal/mtlrendercom...

magic_at_enimai 5y ago

Here are the matmul sizes for the MiniLM model used for inference: https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/bert_minilm_l12_h384_matmul.txt https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/b... These are the matmul sizes for the BERT training workload https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/benchmark_small_sizes.txt https://github.com/mmperf/mmperf/blob/main/benchmark_sizes/b... Yes we use the latest MoltenVK (1.3.204.0) installed in the system. I will let @noxa and other IREE devs chime in on the SPIR-V path but we do support prefix sums etc in the GPU path. //part of nod.ai team.

raphlinus 5y ago
Thanks for the matmul sizes, but the question I am more interested in is precision. Matrix multiply throughput can be dramatically higher for tensor cores[1] than normal shader ALU, specifically for reduced precision arithmetic. I'm wondering to what extent that's accessible on the M1, and to what extent IREE can address them. Regarding prefix sum, the specific question I'm interested in is that SPIRV OpControlBarrier with device scope gets translated into threadgroup_barrier(mem_device) [2]. That's insufficient to make decoupled look-back work. Conversely, if you're not using decoupled look-back, you're not getting the full throughput on GPUs that do support that barrier. I'm wondering how your infrastructure deals with that. [1]: https://developer.nvidia.com/blog/programming-tensor-cores-cuda-9/ https://developer.nvidia.com/blog/programming-tensor-cores-c... [2]: https://github.com/linebender/piet-gpu/blob/d81e5cb4ee145abd1f2f01904579cb3f93960e97/tests/shader/gen/prefix.msl#L136 https://github.com/linebender/piet-gpu/blob/d81e5cb4ee145abd...
- magic_at_enimai 5y ago
  So with Tensorcores you use TF32 which is more like FP19-ish and the marketing makes you think you get 8x the performance. But if you want actual FP32 precision you will need something like [1] but then your performance in the Tensorcore path is _only_ 2X faster than the SIMT path. I'll leave the prefix sum for other devs who know more :D https://github.com/NVIDIA/cutlass/blob/master/examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm/27_ampere_3xtf32_fast_accurate_tensorop_gemm.cu https://github.com/NVIDIA/cutlass/blob/master/examples/27_am... //part of nod.ai/shark team
  - raphlinus 5y ago
    I think we're talking past each other to some extent. Putting aside the question of how misleading it is to market a 16 bit multiply as a "TF32" operation, this is all about tradeoffs. The specific tradeoff that these tensor cores make is that in exchange for reduced precision (and a programming model which is even more of a pain than ordinary compute shaders, an astonishing achievement in and of itself), you get a lot more throughput. For certain AI workloads, particularly inference, that tradeoff is well worth it. Reading between the lines a little, it sounds like your infrastructure is potentially able to exploit a good deal of the available throughput for FP32 workloads. That's great, and I'm happy to see it! However, for workloads that don't need that much precision, the tradeoff might be a lot less advantageous to M1. That may change again if and when Apple opens up lower-level APIs to their hardware, or reverse engineering delivers usable results.
    - shaklee3 5y ago
      tf32 and fp16 tensor cores are completely different, and tf32 is not 16 bit multiplication.

noxa 5y ago

RE cooperative matrix: they have functional units (AMX/ANE) that could hopefully be exposed in MSL via something shaped like cooperative matrix and I'm pretty sure it'd be fantastic. Everything today is locked behind CoreML and Accelerate and those are poor targets for modern compiler-based approaches :( On the Vulkan side there's been rumblings of a vendor-agnostic extension for cooperative matrix and support from major vendors - at which point I'm hoping that leads to Apple wanting to show off their own HW features. RE MoltenVK: we were surprised how robust it is nowadays - we're definitely going to build out our own Metal backend for our abstraction layer but wanted to see what we could hit with the zero-code option and it's proven very useful for that! A MoltenVK build is on the order of ~12MB (last I looked) while the entire IREE runtime is 50-100KB so it's a hard pill to swallow, all other issues (security/memory consumption/startup time/etc) notwithstanding :) AFAIK the memoryless storage is only for textures and mostly useful in render passes - Vulkan has this via VK_MEMORY_PROPERTY_LAZILY_ALLOCATED_BIT. In a way what we do in our compiler is fuse dispatches such that they never hit memory at all so it hasn't been something we've needed yet. For heavy compute workloads more control over virtual memory would be useful, though, ala https://developer.nvidia.com/blog/introducing-low-level-gpu-virtual-memory-management/ https://developer.nvidia.com/blog/introducing-low-level-gpu-... - some of the biggest performance hurdles when it comes to highly dynamic ML is around memory management (being able to zero-copy resize buffers when data-dependent shapes change would be killer). // IREE dev

raphlinus 5y ago
Great, thanks. That answers my questions. I'll read up on the lazily allocated bit; I wasn't aware that this provided similar functionality as dispatchThreadsPerTile[1], but perhaps it's something I'm misunderstanding. I'm excited about that as a way to stitch 2D graphics rendering operations together without having to hit main memory, but from your explanation I can see that functionality might not be very useful in AI workloads. Amen on more control over dynamic memory access patterns. It's something I'm struggling with too, and I have a feeling that whatever solution I come up with is going to be a compromise. Keep up the good work, these are exciting times! [1]: https://developer.apple.com/documentation/metal/mtlrendercommandencoder/2866171-dispatchthreadspertile https://developer.apple.com/documentation/metal/mtlrendercom...