13 ms·

PyTorch on Apple M1 MAX GPUs with SHARK – faster than TensorFlow-Metal

154 points by yayr 5y ago

yayr 5y ago
there is also some hint about performance/watt vs an NVIDIA A100 using BERT models: > The A100 draws 131W peak during the training run, the M1MAX GPU draws a maximum of 15.4 W. The A100 runs an iteration at 7ms vs 104ms (with SHARK) and 196ms (with TF-Metal)
- sudosysgen 5y ago
  So the A100 is 75% more efficient despite being in a high power variation of a larger node. Seems Apple still haw a lot of work to do in the GPU/ML side of things.
  - olliej 5y ago
    My reading implied that they weren't able to use the apple hardware acceleration for AI, which given the massive perf delta I suspect the A100 is using. If they weren't using the A100's ML hardware I'd be both amazed by Nvidia's engineers, and also unhappy as it would be unfair for the test to not use the features of the hardware (this is a common refrain in some benchmarking circles where you disable features to get "fairness" in comparisons). Honestly though I'm more concerned about the other aspects of apple's software: "The OSX Window Server crashes when all GPUs are used to the maximum"
    - noxa 5y ago
      Correct! We're currently compiling to SPIR-V and then going through MoltenVK/spirv-cross to get to Metal - and our SPIR-V needs quite a bit of work (currently tuned for Mali, which is a very different architecture). The biggest limiter though is as you note Apple doesn't have a way to use the fancy hardware instructions whereas our CUDA version is using the Tensor Cores. The good news (for us) is that this is all effectively -O1 today; there's still potential for 2-8x speedups over these numbers and that's not even factoring in the inaccessible HW features that maybe one day Apple will expose :crossed-fingers: :) // IREE dev
      - AceJohnny2 5y ago
        Any chance of using the Neural Engine (otherwise exposed by CoreML) ?
        magic_at_nodai 5y ago
        unlikely since the interface from ANE is not public and it may change between hardware versions.
  - yayr 5y ago
    yeah, you could also say the performance is only 57% in this particular test for the M1 vs A100. Still not bad for a laptop though with a first gen chip.
    - smoldesu 5y ago
      The M1 is (supposedly) specialized for ML calculations and on a decidedly denser node, though. I guess it's a testament to either CUDA's power or machine learning GPGPU optimization.
      - yayr 5y ago
        I like this blog to understand NVidia perf stats (search for "Normalized GPU deep learning performance"): https://timdettmers.com/2020/09/07/which-gpu-for-deep-learning/ https://timdettmers.com/2020/09/07/which-gpu-for-deep-learni... It seems that their consumer GPUs also have a gap
      - simondotau 5y ago
        > The M1 is (supposedly) specialized for ML calculations It is, but they’re not using any of that. Only the M1 GPU at this time.
        magic_at_nodai 5y ago
        Yeah the ANE and AMX on cpu are wrapped behind Accelerate Framework and CoreML. So you will have to use CoreML (which wasn't able to compile the latest TF BERT). ANE is also inference only. So if you want training you will have to use the GPU with Apple's Tensorflow-Metal. //part of nod.ai / SHARK team.
deleted 5y ago
[deleted]
phyalow 5y ago
This is a big deal. Congrats to the team who pulled this off.
- magic_at_nodai 5y ago
  Thanks to: LLVM/MLIR --> For the awesome compiler infrastructure IREE --> For the awesome backend to MLIR SHARK/nod.ai --> For adapting IREE for use on various hardware and fine tuning for target hardware. //part of nod.ai / SHARK team
raphlinus 5y ago
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...
amirhirsch 5y ago
Has anyone else actually installed and demonstrated this? (might be the perfect excuse to buy an M1) "All the code the recreate the tests are here and here. However you will need to install PyTorch torchvision from source since torchvision doesn’t have support for M1 yet. You will also need to build SHARK from the apple-m1-max-support branch from the SHARK repository." So I bet you won't get it running within an hour, possibly 24 hours. I'm testing FFTs on GPU. I got cupy running in jupyter today after a day of pain. The error said: "One method of fixing this is to repeatedly uninstall numpy until none is found" -- I had to manually delete directories and reboot to get cupy working -- now I have very fast fourier tranforms, but the internet says the PyTorch implementation is even better on GPU.... so I have to try it... I created a new venv and tried to install pytorch with GPU support but torch.cuda.is_available() returns False. Why is package management such a disaster?
- magic_at_nodai 5y ago
  This is not part of regular pytorch install. If you can build torch-mlir and SHARK from src you can use it. So hopefully soon we can make pip installable packages but for now the interfaces are in constant development so you will have to build from source. //part of nod.ai / SHARK team.
  - amirhirsch 5y ago
    Thanks for your comment, I'm betting that the M1 example demoed here should be easy to build, and I really do want the M1 for my work. I'm just lamenting how annoying it is to install the packaged versions of PyTorch for GPU since that's what I'm fighting with at this very moment!
    - magic_at_enimai 5y ago
      hear your pain and we really want to make it easy (after we make it work). //part of nod.ai / SHARK team.
liquidmetal 5y ago
How do your benchmarks compare against CoreML?
- magic_at_nodai 5y ago
  I updated the blog with the reference. Basically it crashes to compile the model with https://github.com/NodLabs/shark-samples/blob/main/examples/bert_coreml.log https://github.com/NodLabs/shark-samples/blob/main/examples/.... The coremltools converter is very version specific (like all vendor conversion kits) and still on a version of TF I couldn't get on conda. Also it doesn't allow for training and only FP16 for inference with ANE. All our tests were with FP32. //part of nod.ai/shark team.

154 points by yayr 5y ago

yayr 5y ago

there is also some hint about performance/watt vs an NVIDIA A100 using BERT models: > The A100 draws 131W peak during the training run, the M1MAX GPU draws a maximum of 15.4 W. The A100 runs an iteration at 7ms vs 104ms (with SHARK) and 196ms (with TF-Metal)

sudosysgen 5y ago
So the A100 is 75% more efficient despite being in a high power variation of a larger node. Seems Apple still haw a lot of work to do in the GPU/ML side of things.
- olliej 5y ago
  My reading implied that they weren't able to use the apple hardware acceleration for AI, which given the massive perf delta I suspect the A100 is using. If they weren't using the A100's ML hardware I'd be both amazed by Nvidia's engineers, and also unhappy as it would be unfair for the test to not use the features of the hardware (this is a common refrain in some benchmarking circles where you disable features to get "fairness" in comparisons). Honestly though I'm more concerned about the other aspects of apple's software: "The OSX Window Server crashes when all GPUs are used to the maximum"
  - noxa 5y ago
    Correct! We're currently compiling to SPIR-V and then going through MoltenVK/spirv-cross to get to Metal - and our SPIR-V needs quite a bit of work (currently tuned for Mali, which is a very different architecture). The biggest limiter though is as you note Apple doesn't have a way to use the fancy hardware instructions whereas our CUDA version is using the Tensor Cores. The good news (for us) is that this is all effectively -O1 today; there's still potential for 2-8x speedups over these numbers and that's not even factoring in the inaccessible HW features that maybe one day Apple will expose :crossed-fingers: :) // IREE dev
    - AceJohnny2 5y ago
      Any chance of using the Neural Engine (otherwise exposed by CoreML) ?
      - magic_at_nodai 5y ago
        unlikely since the interface from ANE is not public and it may change between hardware versions.
- yayr 5y ago
  yeah, you could also say the performance is only 57% in this particular test for the M1 vs A100. Still not bad for a laptop though with a first gen chip.
  - smoldesu 5y ago
    The M1 is (supposedly) specialized for ML calculations and on a decidedly denser node, though. I guess it's a testament to either CUDA's power or machine learning GPGPU optimization.
    - yayr 5y ago
      I like this blog to understand NVidia perf stats (search for "Normalized GPU deep learning performance"): https://timdettmers.com/2020/09/07/which-gpu-for-deep-learning/ https://timdettmers.com/2020/09/07/which-gpu-for-deep-learni... It seems that their consumer GPUs also have a gap
    - simondotau 5y ago
      > The M1 is (supposedly) specialized for ML calculations It is, but they’re not using any of that. Only the M1 GPU at this time.
      - magic_at_nodai 5y ago
        Yeah the ANE and AMX on cpu are wrapped behind Accelerate Framework and CoreML. So you will have to use CoreML (which wasn't able to compile the latest TF BERT). ANE is also inference only. So if you want training you will have to use the GPU with Apple's Tensorflow-Metal. //part of nod.ai / SHARK team.

deleted 5y ago

[deleted]

phyalow 5y ago

This is a big deal. Congrats to the team who pulled this off.

magic_at_nodai 5y ago
Thanks to: LLVM/MLIR --> For the awesome compiler infrastructure IREE --> For the awesome backend to MLIR SHARK/nod.ai --> For adapting IREE for use on various hardware and fine tuning for target hardware. //part of nod.ai / SHARK team

raphlinus 5y ago

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...

amirhirsch 5y ago

Has anyone else actually installed and demonstrated this? (might be the perfect excuse to buy an M1) "All the code the recreate the tests are here and here. However you will need to install PyTorch torchvision from source since torchvision doesn’t have support for M1 yet. You will also need to build SHARK from the apple-m1-max-support branch from the SHARK repository." So I bet you won't get it running within an hour, possibly 24 hours. I'm testing FFTs on GPU. I got cupy running in jupyter today after a day of pain. The error said: "One method of fixing this is to repeatedly uninstall numpy until none is found" -- I had to manually delete directories and reboot to get cupy working -- now I have very fast fourier tranforms, but the internet says the PyTorch implementation is even better on GPU.... so I have to try it... I created a new venv and tried to install pytorch with GPU support but torch.cuda.is_available() returns False. Why is package management such a disaster?

magic_at_nodai 5y ago
This is not part of regular pytorch install. If you can build torch-mlir and SHARK from src you can use it. So hopefully soon we can make pip installable packages but for now the interfaces are in constant development so you will have to build from source. //part of nod.ai / SHARK team.
- amirhirsch 5y ago
  Thanks for your comment, I'm betting that the M1 example demoed here should be easy to build, and I really do want the M1 for my work. I'm just lamenting how annoying it is to install the packaged versions of PyTorch for GPU since that's what I'm fighting with at this very moment!
  - magic_at_enimai 5y ago
    hear your pain and we really want to make it easy (after we make it work). //part of nod.ai / SHARK team.

liquidmetal 5y ago

How do your benchmarks compare against CoreML?

magic_at_nodai 5y ago
I updated the blog with the reference. Basically it crashes to compile the model with https://github.com/NodLabs/shark-samples/blob/main/examples/bert_coreml.log https://github.com/NodLabs/shark-samples/blob/main/examples/.... The coremltools converter is very version specific (like all vendor conversion kits) and still on a version of TF I couldn't get on conda. Also it doesn't allow for training and only FP16 for inference with ANE. All our tests were with FP32. //part of nod.ai/shark team.