7 ms·

AMD MI300x GPUs with GEMM tuning improves throughput and latency by up to 7.2x

66 points by latchkey 2y ago

teaearlgraycold 2y ago
I get the impression Nvidia thinks they have a moat with CUDA, but the current AI boom is mostly built upon Python libraries that are platform agnostic, or at least can be. With enough support for AMD in Pytorch etc. the decision to buy AMD or Nvidia will purely be down to specs.
- progbits 2y ago
  Almost everyone is using stuff like tensorrt which is far from platform agnostic. But if AMD can build an API-compatible library that performs well it should indeed be easy enough to switch.
  - qeternity 2y ago
    vLLM is widely used (and what is used in this benchmark) which is far easier to setup and maintain than TensorRT. vLLM is not quite as performant but it's pretty close in production environments.
    - KaoruAoiShiho 2y ago
      Is it really that close in production? The benchmarks I've seen show it's almost 3x slower compared to the SOTA. https://bentoml.com/blog/benchmarking-llm-inference-backends https://bentoml.com/blog/benchmarking-llm-inference-backends
      - qeternity 2y ago
        You say benchmarks (plural) but I see just one here, and this benchmark does not add up in my experiences. But let's take L3 8b in fp16 (the first graph) which shows the performance to be comparable on throughput (vLLM faster on TTFT across the board). On the 4bit L3 70b test, they are using AWQ. GPTQ with Marlin kernels (which now gets repacked on the fly in vLLM) is much faster in our tests, and has much better optimizations amongst vLLM contributors. So yeah, this is sample size of n=1, as are my own tests. But my experiences have been much closer to the L3 8b fp16 charts, so I'm going to presume the bigger delta comes down to that. EDIT: here is one of our own recent internal tests where we observed Marlin performing 25-30% better than AWQ in vLLM - https://miro.medium.com/v2/resize:fit:1400/format:webp/1*F9i8w_ytuKO3_BrThfGvqg.png https://miro.medium.com/v2/resize:fit:1400/format:webp/1*F9i...
        KaoruAoiShiho 2y ago
        That's cool but a 25-30% improvement wouldn't make up for the 300% advantage LMDeploy and TRT has over vllm on the 70b benchmark. Where are your results for vllm vs lmdeploy and trt?
        qeternity 2y ago
        I'm not suggesting it's down to that. I'm suggesting that anyone running AWQ on vLLM is probably not running it optimally somewhere else. I don't have a chart handy, but in our tests, TRT was ca. 10% faster (but it's much more difficult to setup, you have to convert models, etc). LMDeploy we tested months ago, maybe it's improved now, but they were making fairly wild claims that we didn't observe. Basically, you shouldn't publish a benchmark without providing all of the details. How long did test run? How were requests staggered? What were the input/output token sizes? A lot of these 256 in / 256 out benchmarks are useless. If your system prompt is 4k tokens, prefill becomes a serious issue. How performant is your continuous batching? Can you do chunked prefill? Are you running prompt prefix caching (and if so, how performant is that)? It's a lot more complex than simply generating some tokens at various batch sizes.
        KaoruAoiShiho 2y ago
        Yeah, that makes sense. Here's another benchmark that shows trt being about 2-3x faster, 2x on fp16 and 3x on int4. https://blog.premai.io/prem-benchmarks/ https://blog.premai.io/prem-benchmarks/ Though to be fair it's 512 tokens and not 4k tokens. Would love to see more public testing done for all the different variations.
        qeternity 2y ago
        Yep, agreed on the more testing. I don't actually have any allegiance to vLLM. We've just been happy with it. A few things in the linked article that make my eyebrows raise. Notably, they claim to achieve 206 t/s at bs=1 on a single A100 80GB (2 TB/s bandwidth) in fp16. That model is 14.5GB in size, and best case would require aggregate memory bandwidth of almost 3 TB/s to achieve. The BentML you linked to, also ran on an A100 80GB, and achieved ~650 t/s at bs=10 so roughly 65 t/s per stream. Granted this is for an 8B param model, not 7B, and it will of course be faster at bs=1 but absolutely not by this margin. In this same benchmark, they achieved roughly 45 t/s per stream at bs=50 so you can see the sort of scaling we're dealing with. At bs=10 you should achieve relatively similar per-stream throughput to bs=10. Baseten did a benchmark with TensorRT-LLM and an A100 80GB on Mistral 7B and at bs=1 they achieved 75 t/s and at bs=8 they achieved 66 t/s per stream (https://www.baseten.co/blog/unlocking-the-full-power-of-nvidia-h100-gpus-for-ml-inference-with-tensorrt/ https://www.baseten.co/blog/unlocking-the-full-power-of-nvid...) I strongly suspect there is something very wrong with the Premai (never heard of them before) or there has been some huge inference breakthrough that I am unaware of. The BentoML benchmarks look pretty good, and I suspect there is some vLLM performance left on the table, however not enough to close the gap. In our testing, TensorRT-LLM was definitely faster, but not enough to warrant all of the other headaches.
- jpgvm 2y ago
  Nvidia should be very afraid of AMD. AMDs mastery of chiplet designs is in stark contrast to Nvidias reliance on extremely large and complex dies and pretty lack luster forays into multi-die in the past. If Nvidia doesn't come up with ways to combat the inherent cost and yield advantages of chiplets they could get quite easily blown away on $/perf. GPGPU didn't have enough money behind it to justify fully exploiting all hardware vs just going with whatever was easiest (Nvidia+CUDA) and eating the cost but that has changed substantially. If you can buy AMD at a 10%+ $/perf advantage when you are talking about laying down $100M+ for accelerators on even a moderate sized cluster you can be damn sure the software will get tuned to take advantage of it.
  - Rinzler89 2y ago
    What are you talking about?! Nvidia has been using chiplets and silicone interconnects in the data center ML accelerators way before AMD got in the game. It's the same TSMC tech that both Nvidia and AMD use, it's not something proprietary exclusive to one or another since they're both fables. On the contrary, AMD has to be afraid since they have no answer to Nvidia's 800GB/s data center interconnects they got from their aquisition of Mellanox, which enables Nvidia to scale.
    - timschmidt 2y ago
      AMD purchased Xilinx not long ago, and as a result has access to their high speed interconnects, including at least one capable of supporting 800G networking: https://www.xilinx.com/products/technology/high-speed-serial.html https://www.xilinx.com/products/technology/high-speed-serial...
    - jpgvm 2y ago
      I'm as close as you get to an old school Infiniband fanboi. Don't think that Infiniband is somehow an Nvidia moat, if anything because they ate up the last somewhat independent vendor (Mellanox) it's likely it will be replaced by some variant of what is now being called Ultra Ethernet which for the most part incorporates ideas from Infiniband into the Ethernet PHY layer along with bolting into libfabric at the userland level. The people with the brains behind it include the Omni-Path folk (which is descendent from QLogic and SilverStorm before that) which have extensive experience in Infiniband and are now free from the shackles of Intel. However this has absolutely nothing to do with chiplet designs or multi-die architectures (which Nvidia has done before with their dual GPU stuff from the GTX 900 series days etc). The fact of the matter is that Nvidia has no experience whatsoever with chiplets or advanced packaging. AMD has on the other hand absolutely tons from yes, chiplets but also HBM, 3d-vcache etc all of which are ultimately advanced packaging techniques that all build on shared background of expertise. Nvidia needs to rapidly catch up in this area or their future is a lot less certain. This is putting aside what is happening with Tenstorrent/Etched/Groq, etc which threaten to eat up the vast majority of the non-training workloads associated with the current crop of AI/LLM architectures.
- almostgotcaught 2y ago
  Why do people make these kinds of comments that dramatically oversimplify incredibly complex issues? Is it some kind of fun "I'm in the in-crowd" type feeling? > mostly built upon Python libraries that are platform agnostic, or at least can be. With enough support for AMD in Pytorch etc You can't imagine how much and how long the work in "can be" is. Thus NVIDIA has two of the deepest moats you can have: time and money (you would probably have to outspend them by 10x for 10 years to catch up). Source: go look at how the AMD backend is implemented in pytorch vs how the AMD backend is implemented. Edit: people must really not understand that this code (ie building this stuff) is highly non-trivial. It's not like building a competing js frontend framework. You not only need good engineers you need lots of them because the surface area is very big - hardware platforms, especially AMD's aren't at all "abstract", so really to do it right you need a team per box in the product matrix.
  - Certhas 2y ago
    I don't see it. Large customers are choosing AMD over nVidia. NVidia has a massive headstart, but there are diminishing returns to their software work. If AMD performance per dollar is 10% better at the hardware level, and their software performance is 90% of that of nvidia they are ahead.
    - WanderPanda 2y ago
      Plus all the flexibility of CUDA is not needed for LLMs anyways
    - almostgotcaught 2y ago
      Absolutely none of this comment is based on fact or reason. > I don't see it. Large customers are choosing AMD over nVidia. Please show me term sheets of contracts signed with "large customers". Please include net, gross, and recurring revenue. > diminishing returns to their software work Makes zero sense: software penetration follows power laws (network effects etc) not exponential decay laws. > If AMD performance per dollar is 10% better at the hardware level, and their software performance is 90% of that of nvidia they are ahead. Again: you cannot fathom how enormous that "if" there is so I highly recommend you actually go look at the impl to get a sense.
- deleted 2y ago
  [deleted]
- roenxi 2y ago
  There is also the MI300A APU (not the MI300X that the article is about). Do I have any idea about the technical details? No. But IMO CUDA's value prop might actually be memory management and APU seem to be a strategic stab right into the heart of that by making memory management less of an issue. If the GPU and CPU start sharing the same memory, a lot of the value proposition of CUDA goes away. Haven't seen that land meaningfully so far, but it is an interesting risk to Nvidia.
  - ColonelPhantom 2y ago
    What makes you think that CUDA's value prop is memory management? I have seen no indication that implementing the memory allocation system is the hard part of any GPGPU system, and HIP is mostly source-compatible with CUDA, so it's not like the user has to do anything different memory management-wise to run on AMD.
    - roenxi 2y ago
      I owned an AMD card. Everything high-level I wanted to do I could implement on the CPU fine, but I couldn't implement it on the GPU because it was too hard to operate on data in VRAM without crashing (been a while, I forget if it was the transfer or the operation that failed), and to get back good debug info because I needed to move data from VRAM to normal RAM. I was trying to learn at the time so efficiency wasn't really a factor, I just wanted things to work. They never did. CPU was easy, GPU was hard. The major difference I could spot though was all the time I was spending batting data around between different buffers, the APIs weren't really limiting me.
      - kllrnohj 2y ago
        Getting good performance out of the GPU in general is regularly a memory problem. CUDA doesn't really help beyond documenting what is expected of you with regards to data layout and warp coalescing. https://docs.nvidia.com/cuda/cuda-c-programming-guide/#maximize-memory-throughput https://docs.nvidia.com/cuda/cuda-c-programming-guide/#maxim...
- llm_trw 2y ago
  This is what everyone who hasn't spend months fighting AMD drivers thinks ... until they spend months fighting AMD drivers. Geohot's streams show you why that's a bad idea in real time.
  - Workaccount2 2y ago
    AMD's ability to not put out good software for 20 years now is probably one of the most impressive feats in all of computing history.
    - llm_trw 2y ago
      The most shocking thing AMD's done in the last 20 years is not fucking up Ryzen after the first release. I have no idea how they managed that give their corporate DNA.
      - lostmsu 2y ago
        Anecdotally mine is unstable.
- threatripper 2y ago
  Running inference with LLMs is not very complex. You don't need CUDA for that at all. You don't even need those complex libraries for inference. The market for inference will grow much more as AI is used in more and more applications by more and more people. NVIDIA has a moat but it's around the smaller castle.
  - YetAnotherNick 2y ago
    You need CUDA for flash attention. Other than that its just matrix multiplication, embedding and point wise activation.
    - f_devd 2y ago
      Only for the reference implementation, flash attention is just based on optimization of the memory bandwidth between caches. The ThunderKittens[0] library has what is effectively flash attention v3, and they are working on supporting AMD arch. [0]: https://hazyresearch.stanford.edu/blog/2024-05-12-tk https://hazyresearch.stanford.edu/blog/2024-05-12-tk
- nabla9 2y ago
  CUDA and ROCm work under Pytorch. If ROCm does not work well Pytorch does not work well. Nvidia has multiple advantages. 1. Same software written in PyTorch works across all relatively modern Nvidia chips. With AMD and ROCm that's not the case. 2. M̶I̶3̶0̶0̶x̶ h̶a̶s̶ n̶o̶ F̶P̶8̶ s̶u̶p̶p̶o̶r̶t̶. 3. Comparisons against H100 are always like this: 8x AMD MI300X (192GB, 750W) GPU 8x H100 SXM5 (80GB, 700W) GPU The fair comparison would be against 8x H100 NVL (188GB, <800W) GPU And that they never do. 4. H100 is 21 month old architecture. MI300X is 7 months. Nvidia moved into new architecture every year pace. AMD is generation behind and must step up the pace. B100 comes out this year. AMD is getting closer, but don't except them to catch Nvidia in not time.
  - Palmik 2y ago
    It has fp8 support. Not sure whether fp8 on MI300x is supported by vLLM yet. Also, many of these comparisons use vLLM for both setups, but for Nvidia you can and should use TensorRT-LLM which tends to do quite a bit better than vLLM at high loads.
    - deleted 2y ago
      [deleted]
    - latchkey 2y ago
      Elio, the person who did the testing confirmed with me that he has fp8 working.
  - latchkey 2y ago
    > MI300X is 7 months. Less than that, we paid for ours in January and received it in March. The first batch had problems and we had to send them back, which took another 3 weeks. So, let's consider the start date closer to April. ~3 months.
  - cromulen 2y ago
    I haven't ever seen a server with 8x H100 NVL 188GB. The H100 NVL has 94GB of VRAM but they sell them in pairs connected with NVLink, so I guess they sometimes market them as 188GB, but in fact it's two cards and a server usually has 4 pairs.
- drexlspivey 2y ago
  AMD has sold out their MI300x and are trying to get more allocation from TSMC. NVDA's moat is their relationship with TSMC.
  - latchkey 2y ago
    Eh? I just ordered 128 of them a few days ago.
    - paulmd 2y ago
      in fairness, you also are a high-profile cloud provider with a pre-existing relationship - if small. ;)
      - latchkey 2y ago
        Are you having trouble with your allocation?
    - Jlagreen 2y ago
      To be fair, 128 are maybe $2.5 million revenue for AMD. If they predict $4b then we talk about 200-400k MI300X. For every $1b more revenue, AMD needs capacity from TSMC in the area of 50-100k. To be frank, your 128 are a rounding error in that. Nvidia has shipped almost 4 million DC GPUs in 2023 to give your an idea about the sizes for it to count on the balance sheet.
      - latchkey 2y ago
        Mi300x only really started being used in April…. A whopping 3 months ago. We are a brand new startup and tiny. Not pretending to be anything else. But we are also early and one of the first to buy and receive these cards. We have excellent backing and will grow with demand. The more people wake up to an alternative, that is actually available and as we are proving… viable, the more demand we will have. The need for compute is endless. Dell is a partner of ours and sees the value in our proposition. That is huge. They are hungry to grow their AI business the same way AMD is. Give it time. The frankness is appreciated, but obvious. Comparisons against nvidia today are laughable. This is a long tail game. Talk to me in a few years. This is not a football game with a winner and loser. Our goal is to not just be AMD, but every chip that our customers want.
- quelsolaar 2y ago
  If you are buying hardware for billions of dollars, you can afford to spend a fair bit of engineering effort to make your code run on less expencive hardware.
sva_ 2y ago
> Let’s delve into the notable advancements achieved through GEMM tuning of LLMs such as Llama Smells like written by GPT
- robko 2y ago
  I agree. This article is clearly written with a generative language model. A few other telltale signs: 1. Repeating the same thing multiple times with slight variation: * "allowing developers to fine-tune their applications and unlock the full potential of their underlying hardware, ultimately maximising vLLM performance." (fine-tune, unlock potential, maximize performance are all roughly the same thing) * "AI and machine learning models" (AI and machine learning models are the same thing in the context of this article) * "utilise multiple threads or cores" (Why differentiate between threads and cores?) * "tailored to enhance computational efficiency and overall throughput" (efficiency and throughput are highly related) * "a series of graphs and data visualisations" (all the data visualizations in this article are graphs) * "more computational effort and time" (same thing) * "significantly enhanced the performance and efficiency" (same thing) * "ensuring efficient processing and superior performance for complex and demanding AI workloads" (same things) 2. Explaining what "rocBLAS" stands for multiple times. 3. Other ChatGPTisms: * "offering a comprehensive view of [...]" * "Let’s delve into the notable advancements achieved through [...]" * "ensures quicker processing times, which is crucial for [...]" * "effectively mitigated these impacts, maintaining [...]" * "elucidate the impact of" * "significantly enhanced" * "These results underscore the critical role of [...]" * "Key Aspects", "Key Observations", "Key findings" So why is this bad? - Because it undermines the trust in the the article. We do not know whether the claims are actually true or whether they were just made up by ChatGPT.
  - idonotknowwhy 2y ago
    Thanks for articulating something I've noticed happening all over the internet and even in YouTube video scripts. Are claudisms different from gptisms? Why can't these authors tell ChatGPT to write with a different prose and avoid "delve", "crucial", etc? Wizardlm writes these same things in all its answers too
fancyfredbot 2y ago
Can someone knowledgeable please put these numbers in context with a comparison to an H100? For example they say 1053 tokens per second throughput with batch size 4 on llama2 70b. Is that good?
- qeternity 2y ago
  Where do you see 1053 t/s for L2 70b? From their benchmark table, they achieve 1.8s response time at bs=4 for 256 output tokens, which would be 569 t/s throughput. At bs=1 you are memory bandwidth bound. So as a simple ceiling we divide the memory bandwidth by the model size (70b is roughly 140gb at fp16). That would yield roughly 24 t/s for H100 and 39 t/s for MI300x. I'm not really sure how GEMM optimizations alone would yield these sort of improvements at fp16. With quantization, where dequant overheads start to become compute bound, or at very high batch sizes, it makes sense and there has been a lot of work with various kernels in this area. But I'm a bit confused at fp16 and bs=1. Maybe someone smarter can chime in? EDIT: Ok I watched their video, which shows 39.75 t/s in the unoptimized (which checks out given the above rough maths) and 313.95 t/s for the optimized. I'm not really sure what's going on here.
  - fancyfredbot 2y ago
    I found 1053 t/s in cell G18 in the throughput tab of the full results sheet they linked to: https://docs.google.com/spreadsheets/u/0/d/1hHd9riWh-hQJnNOzEI_ZZN1LpdY_bBSbvDOVPdfYslM/htmlview# https://docs.google.com/spreadsheets/u/0/d/1hHd9riWh-hQJnNOz...
    - qeternity 2y ago
      Hmm the post now links to this doc where I got my figures from: https://docs.google.com/spreadsheets/d/1hHd9riWh-hQJnNOzEI_ZZN1LpdY_bBSbvDOVPdfYslM/edit?gid=1853523197#gid=1853523197 https://docs.google.com/spreadsheets/d/1hHd9riWh-hQJnNOzEI_Z... This sheet only lists the response latency, but it does look more recent than the one you listed. In any event, the sheet you link to shows generation speed and throughput (generation speed is cell E18 which shows 527 t/s). G18 is perhaps prefill + gen but even that doesn't add up.
saagarjha 2y ago
I don't usually like to do this, because clearly someone wrote this post (or at the very least they cared enough to have a LLM help draft it :P). Maybe it's just me and my interests not being aligned with the content. But I just found myself kind of disappointed by the content? Like, the post is about how you can do GEMM tuning on AMD GPUs, a subject which is inherently super interesting–there's a lot of nuance to writing optimal kernels, and some of this is expressed in the article, too. Combine that with an architecture that isn't Nvidia? It's an excellent setup for something that would be interesting to read. Which makes the actual conclusion all the more disappointing, IMO. There's nothing about what the actual optimizations are. It's just "oh yeah we tuned our GEMMs and now LLaMA is faster". Like, I get that nobody actually cares about GEMM and they just want tokens to come out of their GPU. But still, that's like writing a blog post about how you can speed up your game with SIMD and then posting some charts of how Cyberpunk 2077 gives you 2x the frame rate now. Ok, but how? I just feel like the interesting part is missing.
- xoranth 2y ago
  I believe the reason they offer no details about how they tuned the kernels is that the tuning is done by a tool provided by AMD. See here: https://rocm.docs.amd.com/projects/rocBLAS/en/develop/how-to/Programmers_Guide.html#rocblas-gemm-tune https://rocm.docs.amd.com/projects/rocBLAS/en/develop/how-to...
- latchkey 2y ago
  This article builds upon the work that C&C just released an article on. https://chipsandcheese.com/2024/06/25/testing-amds-giant-mi300x/ https://chipsandcheese.com/2024/06/25/testing-amds-giant-mi3... Elio, the person who did the testing, worked together with C&C, on my system. He then took it to the next level with this gemm work. The docker container is linked on the article now so that anyone can reproduce the results.
deleted 2y ago
[deleted]
Lindon4290 2y ago
Now, I don't have any MI300X, so I can't make any definite claims here. I am hoping someone else can replicate the results shown here or at the least educate me on how this is possible. Good part is the docker container and associated steps are made public - which is pretty cool! Going by the video, the first thing that gave me pause was that a single MI300X is pulling off groq like performance, i.e. 314 tokens/second for a batch size = 1 (bs=1) with a prompt of 256 tokens and generation of 256 tokens. [1] The Llama-2 70B is 128.48GB with FP16 (you can see this in the video). The entire model fits well within the 192GB HBM memory of the MI300X - which is awesome! However, for an regressive transformer model, during generation, the entire model weights are processed to generate a single next token. These models are "next token predictors" so to say, and you need the previous token to generate the next token. Therefore, the 128.48GB of model weights need are consumed from the HBM at the compute cores of the MI300X, per generated token. Note, I am not talking about the prefill - which only needs a single forward pass to generate the first output token. Every subsequent output token is auto-regressive. The video shows that a single prompt (bs=1) with 256 token prompt and 256 tokens generated within 1.63 second. There is no tensor parallelism involved, or batching or anything else. This is a bs=1 case with a single card, and hence you can reason about the math fairly easy. This shouldn't fundamentally be possible within the specs of the MI300X card. The card has a peak HBM memory bandwidth of 5.3 TB/s. You'll notice that to cycle through the weights (assuming FP16) 256 times, you'd need a minimum of 6 seconds, even at perfect ideal conditions. Napkin math: (256 * 128.48e9) / (5.3e12) [1] https://wow.groq.com/groq-sets-new-large-language-model-performance-record-of-300-tokens-per-second-per-user-on-meta-ai-foundational-llm-llama-2-70b/ https://wow.groq.com/groq-sets-new-large-language-model-perf...
- latchkey 2y ago
  Elio clarified this to me, he said... 256 tokens in 1 prompt/batch Not 256 batches
  - Lindon4290 2y ago
    Yes, a single sequence with 256 prompt tokens and 256 output tokens. This is a batch size = 1. No one is saying anything about 256 batches. The first step in understanding this is to notice that the model (llama2) generates 1 output token at a time. This is because the llama2 70B is a autoregressive decoder-only transformer. Fundamentally, to generate a single output token you need to process the entire model weights. At each forward pass you generate 1 token. OK, now to generate 256 output tokens - you need 256 sequential forward passes. At each forward pass, the entire model is read from the gpu VRAM. Even at ideal memory bandwidth (5.3 TB/s) that (256 forward passes of a 128.48GB model) should take 6s. The reported number of 1.63s should not be possible. I'd strongly recommend checking for correctness - that the generate output is coherent. Try sending actual prompts to the "gemm tuned" model and observing the generated responses and latencies. With the "benchmark_throughput.py" you only get a final number and there is no check whether the output is valid or not.
fancyfredbot 2y ago
The numbers in this article don't make sense. They aren't consistent with the hardware (they seem to show the weights being loaded faster than mi300x peak memory bandwidth) and they aren't self consistent (70B models running only 2x slower than 7B models). I'm not sure what's going on here but something is not right.
- Lindon4290 2y ago
  Yeah, I've asked questions along those lines as well. Something sketchy is going on. See here - https://news.ycombinator.com/item?id=40833109 https://news.ycombinator.com/item?id=40833109
  - deleted 2y ago
    [deleted]

66 points by latchkey 2y ago

I get the impression Nvidia thinks they have a moat with CUDA, but the current AI boom is mostly built upon Python libraries that are platform agnostic, or at least can be. With enough support for AMD in Pytorch etc. the decision to buy AMD or Nvidia will purely be down to specs.

progbits 2y ago
Almost everyone is using stuff like tensorrt which is far from platform agnostic. But if AMD can build an API-compatible library that performs well it should indeed be easy enough to switch.
- qeternity 2y ago
  vLLM is widely used (and what is used in this benchmark) which is far easier to setup and maintain than TensorRT. vLLM is not quite as performant but it's pretty close in production environments.
  - KaoruAoiShiho 2y ago
    Is it really that close in production? The benchmarks I've seen show it's almost 3x slower compared to the SOTA. https://bentoml.com/blog/benchmarking-llm-inference-backends https://bentoml.com/blog/benchmarking-llm-inference-backends
    - qeternity 2y ago
      You say benchmarks (plural) but I see just one here, and this benchmark does not add up in my experiences. But let's take L3 8b in fp16 (the first graph) which shows the performance to be comparable on throughput (vLLM faster on TTFT across the board). On the 4bit L3 70b test, they are using AWQ. GPTQ with Marlin kernels (which now gets repacked on the fly in vLLM) is much faster in our tests, and has much better optimizations amongst vLLM contributors. So yeah, this is sample size of n=1, as are my own tests. But my experiences have been much closer to the L3 8b fp16 charts, so I'm going to presume the bigger delta comes down to that. EDIT: here is one of our own recent internal tests where we observed Marlin performing 25-30% better than AWQ in vLLM - https://miro.medium.com/v2/resize:fit:1400/format:webp/1*F9i8w_ytuKO3_BrThfGvqg.png https://miro.medium.com/v2/resize:fit:1400/format:webp/1*F9i...
      - KaoruAoiShiho 2y ago
        That's cool but a 25-30% improvement wouldn't make up for the 300% advantage LMDeploy and TRT has over vllm on the 70b benchmark. Where are your results for vllm vs lmdeploy and trt?
        qeternity 2y ago
        I'm not suggesting it's down to that. I'm suggesting that anyone running AWQ on vLLM is probably not running it optimally somewhere else. I don't have a chart handy, but in our tests, TRT was ca. 10% faster (but it's much more difficult to setup, you have to convert models, etc). LMDeploy we tested months ago, maybe it's improved now, but they were making fairly wild claims that we didn't observe. Basically, you shouldn't publish a benchmark without providing all of the details. How long did test run? How were requests staggered? What were the input/output token sizes? A lot of these 256 in / 256 out benchmarks are useless. If your system prompt is 4k tokens, prefill becomes a serious issue. How performant is your continuous batching? Can you do chunked prefill? Are you running prompt prefix caching (and if so, how performant is that)? It's a lot more complex than simply generating some tokens at various batch sizes.
        KaoruAoiShiho 2y ago
        Yeah, that makes sense. Here's another benchmark that shows trt being about 2-3x faster, 2x on fp16 and 3x on int4. https://blog.premai.io/prem-benchmarks/ https://blog.premai.io/prem-benchmarks/ Though to be fair it's 512 tokens and not 4k tokens. Would love to see more public testing done for all the different variations.
        qeternity 2y ago
        Yep, agreed on the more testing. I don't actually have any allegiance to vLLM. We've just been happy with it. A few things in the linked article that make my eyebrows raise. Notably, they claim to achieve 206 t/s at bs=1 on a single A100 80GB (2 TB/s bandwidth) in fp16. That model is 14.5GB in size, and best case would require aggregate memory bandwidth of almost 3 TB/s to achieve. The BentML you linked to, also ran on an A100 80GB, and achieved ~650 t/s at bs=10 so roughly 65 t/s per stream. Granted this is for an 8B param model, not 7B, and it will of course be faster at bs=1 but absolutely not by this margin. In this same benchmark, they achieved roughly 45 t/s per stream at bs=50 so you can see the sort of scaling we're dealing with. At bs=10 you should achieve relatively similar per-stream throughput to bs=10. Baseten did a benchmark with TensorRT-LLM and an A100 80GB on Mistral 7B and at bs=1 they achieved 75 t/s and at bs=8 they achieved 66 t/s per stream (https://www.baseten.co/blog/unlocking-the-full-power-of-nvidia-h100-gpus-for-ml-inference-with-tensorrt/ https://www.baseten.co/blog/unlocking-the-full-power-of-nvid...) I strongly suspect there is something very wrong with the Premai (never heard of them before) or there has been some huge inference breakthrough that I am unaware of. The BentoML benchmarks look pretty good, and I suspect there is some vLLM performance left on the table, however not enough to close the gap. In our testing, TensorRT-LLM was definitely faster, but not enough to warrant all of the other headaches.
jpgvm 2y ago
Nvidia should be very afraid of AMD. AMDs mastery of chiplet designs is in stark contrast to Nvidias reliance on extremely large and complex dies and pretty lack luster forays into multi-die in the past. If Nvidia doesn't come up with ways to combat the inherent cost and yield advantages of chiplets they could get quite easily blown away on $/perf. GPGPU didn't have enough money behind it to justify fully exploiting all hardware vs just going with whatever was easiest (Nvidia+CUDA) and eating the cost but that has changed substantially. If you can buy AMD at a 10%+ $/perf advantage when you are talking about laying down $100M+ for accelerators on even a moderate sized cluster you can be damn sure the software will get tuned to take advantage of it.
- Rinzler89 2y ago
  What are you talking about?! Nvidia has been using chiplets and silicone interconnects in the data center ML accelerators way before AMD got in the game. It's the same TSMC tech that both Nvidia and AMD use, it's not something proprietary exclusive to one or another since they're both fables. On the contrary, AMD has to be afraid since they have no answer to Nvidia's 800GB/s data center interconnects they got from their aquisition of Mellanox, which enables Nvidia to scale.
  - timschmidt 2y ago
    AMD purchased Xilinx not long ago, and as a result has access to their high speed interconnects, including at least one capable of supporting 800G networking: https://www.xilinx.com/products/technology/high-speed-serial.html https://www.xilinx.com/products/technology/high-speed-serial...
  - jpgvm 2y ago
    I'm as close as you get to an old school Infiniband fanboi. Don't think that Infiniband is somehow an Nvidia moat, if anything because they ate up the last somewhat independent vendor (Mellanox) it's likely it will be replaced by some variant of what is now being called Ultra Ethernet which for the most part incorporates ideas from Infiniband into the Ethernet PHY layer along with bolting into libfabric at the userland level. The people with the brains behind it include the Omni-Path folk (which is descendent from QLogic and SilverStorm before that) which have extensive experience in Infiniband and are now free from the shackles of Intel. However this has absolutely nothing to do with chiplet designs or multi-die architectures (which Nvidia has done before with their dual GPU stuff from the GTX 900 series days etc). The fact of the matter is that Nvidia has no experience whatsoever with chiplets or advanced packaging. AMD has on the other hand absolutely tons from yes, chiplets but also HBM, 3d-vcache etc all of which are ultimately advanced packaging techniques that all build on shared background of expertise. Nvidia needs to rapidly catch up in this area or their future is a lot less certain. This is putting aside what is happening with Tenstorrent/Etched/Groq, etc which threaten to eat up the vast majority of the non-training workloads associated with the current crop of AI/LLM architectures.
almostgotcaught 2y ago
Why do people make these kinds of comments that dramatically oversimplify incredibly complex issues? Is it some kind of fun "I'm in the in-crowd" type feeling? > mostly built upon Python libraries that are platform agnostic, or at least can be. With enough support for AMD in Pytorch etc You can't imagine how much and how long the work in "can be" is. Thus NVIDIA has two of the deepest moats you can have: time and money (you would probably have to outspend them by 10x for 10 years to catch up). Source: go look at how the AMD backend is implemented in pytorch vs how the AMD backend is implemented. Edit: people must really not understand that this code (ie building this stuff) is highly non-trivial. It's not like building a competing js frontend framework. You not only need good engineers you need lots of them because the surface area is very big - hardware platforms, especially AMD's aren't at all "abstract", so really to do it right you need a team per box in the product matrix.
- Certhas 2y ago
  I don't see it. Large customers are choosing AMD over nVidia. NVidia has a massive headstart, but there are diminishing returns to their software work. If AMD performance per dollar is 10% better at the hardware level, and their software performance is 90% of that of nvidia they are ahead.
  - WanderPanda 2y ago
    Plus all the flexibility of CUDA is not needed for LLMs anyways
  - almostgotcaught 2y ago
    Absolutely none of this comment is based on fact or reason. > I don't see it. Large customers are choosing AMD over nVidia. Please show me term sheets of contracts signed with "large customers". Please include net, gross, and recurring revenue. > diminishing returns to their software work Makes zero sense: software penetration follows power laws (network effects etc) not exponential decay laws. > If AMD performance per dollar is 10% better at the hardware level, and their software performance is 90% of that of nvidia they are ahead. Again: you cannot fathom how enormous that "if" there is so I highly recommend you actually go look at the impl to get a sense.
deleted 2y ago
[deleted]
roenxi 2y ago
There is also the MI300A APU (not the MI300X that the article is about). Do I have any idea about the technical details? No. But IMO CUDA's value prop might actually be memory management and APU seem to be a strategic stab right into the heart of that by making memory management less of an issue. If the GPU and CPU start sharing the same memory, a lot of the value proposition of CUDA goes away. Haven't seen that land meaningfully so far, but it is an interesting risk to Nvidia.
- ColonelPhantom 2y ago
  What makes you think that CUDA's value prop is memory management? I have seen no indication that implementing the memory allocation system is the hard part of any GPGPU system, and HIP is mostly source-compatible with CUDA, so it's not like the user has to do anything different memory management-wise to run on AMD.
  - roenxi 2y ago
    I owned an AMD card. Everything high-level I wanted to do I could implement on the CPU fine, but I couldn't implement it on the GPU because it was too hard to operate on data in VRAM without crashing (been a while, I forget if it was the transfer or the operation that failed), and to get back good debug info because I needed to move data from VRAM to normal RAM. I was trying to learn at the time so efficiency wasn't really a factor, I just wanted things to work. They never did. CPU was easy, GPU was hard. The major difference I could spot though was all the time I was spending batting data around between different buffers, the APIs weren't really limiting me.
    - kllrnohj 2y ago
      Getting good performance out of the GPU in general is regularly a memory problem. CUDA doesn't really help beyond documenting what is expected of you with regards to data layout and warp coalescing. https://docs.nvidia.com/cuda/cuda-c-programming-guide/#maximize-memory-throughput https://docs.nvidia.com/cuda/cuda-c-programming-guide/#maxim...
llm_trw 2y ago
This is what everyone who hasn't spend months fighting AMD drivers thinks ... until they spend months fighting AMD drivers. Geohot's streams show you why that's a bad idea in real time.
- Workaccount2 2y ago
  AMD's ability to not put out good software for 20 years now is probably one of the most impressive feats in all of computing history.
  - llm_trw 2y ago
    The most shocking thing AMD's done in the last 20 years is not fucking up Ryzen after the first release. I have no idea how they managed that give their corporate DNA.
    - lostmsu 2y ago
      Anecdotally mine is unstable.
threatripper 2y ago
Running inference with LLMs is not very complex. You don't need CUDA for that at all. You don't even need those complex libraries for inference. The market for inference will grow much more as AI is used in more and more applications by more and more people. NVIDIA has a moat but it's around the smaller castle.
- YetAnotherNick 2y ago
  You need CUDA for flash attention. Other than that its just matrix multiplication, embedding and point wise activation.
  - f_devd 2y ago
    Only for the reference implementation, flash attention is just based on optimization of the memory bandwidth between caches. The ThunderKittens[0] library has what is effectively flash attention v3, and they are working on supporting AMD arch. [0]: https://hazyresearch.stanford.edu/blog/2024-05-12-tk https://hazyresearch.stanford.edu/blog/2024-05-12-tk
nabla9 2y ago
CUDA and ROCm work under Pytorch. If ROCm does not work well Pytorch does not work well. Nvidia has multiple advantages. 1. Same software written in PyTorch works across all relatively modern Nvidia chips. With AMD and ROCm that's not the case. 2. M̶I̶3̶0̶0̶x̶ h̶a̶s̶ n̶o̶ F̶P̶8̶ s̶u̶p̶p̶o̶r̶t̶. 3. Comparisons against H100 are always like this: 8x AMD MI300X (192GB, 750W) GPU 8x H100 SXM5 (80GB, 700W) GPU The fair comparison would be against 8x H100 NVL (188GB, <800W) GPU And that they never do. 4. H100 is 21 month old architecture. MI300X is 7 months. Nvidia moved into new architecture every year pace. AMD is generation behind and must step up the pace. B100 comes out this year. AMD is getting closer, but don't except them to catch Nvidia in not time.
- Palmik 2y ago
  It has fp8 support. Not sure whether fp8 on MI300x is supported by vLLM yet. Also, many of these comparisons use vLLM for both setups, but for Nvidia you can and should use TensorRT-LLM which tends to do quite a bit better than vLLM at high loads.
  - deleted 2y ago
    [deleted]
  - latchkey 2y ago
    Elio, the person who did the testing confirmed with me that he has fp8 working.
- latchkey 2y ago
  > MI300X is 7 months. Less than that, we paid for ours in January and received it in March. The first batch had problems and we had to send them back, which took another 3 weeks. So, let's consider the start date closer to April. ~3 months.
- cromulen 2y ago
  I haven't ever seen a server with 8x H100 NVL 188GB. The H100 NVL has 94GB of VRAM but they sell them in pairs connected with NVLink, so I guess they sometimes market them as 188GB, but in fact it's two cards and a server usually has 4 pairs.
drexlspivey 2y ago
AMD has sold out their MI300x and are trying to get more allocation from TSMC. NVDA's moat is their relationship with TSMC.
- latchkey 2y ago
  Eh? I just ordered 128 of them a few days ago.
  - paulmd 2y ago
    in fairness, you also are a high-profile cloud provider with a pre-existing relationship - if small. ;)
    - latchkey 2y ago
      Are you having trouble with your allocation?
  - Jlagreen 2y ago
    To be fair, 128 are maybe $2.5 million revenue for AMD. If they predict $4b then we talk about 200-400k MI300X. For every $1b more revenue, AMD needs capacity from TSMC in the area of 50-100k. To be frank, your 128 are a rounding error in that. Nvidia has shipped almost 4 million DC GPUs in 2023 to give your an idea about the sizes for it to count on the balance sheet.
    - latchkey 2y ago
      Mi300x only really started being used in April…. A whopping 3 months ago. We are a brand new startup and tiny. Not pretending to be anything else. But we are also early and one of the first to buy and receive these cards. We have excellent backing and will grow with demand. The more people wake up to an alternative, that is actually available and as we are proving… viable, the more demand we will have. The need for compute is endless. Dell is a partner of ours and sees the value in our proposition. That is huge. They are hungry to grow their AI business the same way AMD is. Give it time. The frankness is appreciated, but obvious. Comparisons against nvidia today are laughable. This is a long tail game. Talk to me in a few years. This is not a football game with a winner and loser. Our goal is to not just be AMD, but every chip that our customers want.
quelsolaar 2y ago
If you are buying hardware for billions of dollars, you can afford to spend a fair bit of engineering effort to make your code run on less expencive hardware.

sva_ 2y ago

> Let’s delve into the notable advancements achieved through GEMM tuning of LLMs such as Llama Smells like written by GPT

robko 2y ago
I agree. This article is clearly written with a generative language model. A few other telltale signs: 1. Repeating the same thing multiple times with slight variation: * "allowing developers to fine-tune their applications and unlock the full potential of their underlying hardware, ultimately maximising vLLM performance." (fine-tune, unlock potential, maximize performance are all roughly the same thing) * "AI and machine learning models" (AI and machine learning models are the same thing in the context of this article) * "utilise multiple threads or cores" (Why differentiate between threads and cores?) * "tailored to enhance computational efficiency and overall throughput" (efficiency and throughput are highly related) * "a series of graphs and data visualisations" (all the data visualizations in this article are graphs) * "more computational effort and time" (same thing) * "significantly enhanced the performance and efficiency" (same thing) * "ensuring efficient processing and superior performance for complex and demanding AI workloads" (same things) 2. Explaining what "rocBLAS" stands for multiple times. 3. Other ChatGPTisms: * "offering a comprehensive view of [...]" * "Let’s delve into the notable advancements achieved through [...]" * "ensures quicker processing times, which is crucial for [...]" * "effectively mitigated these impacts, maintaining [...]" * "elucidate the impact of" * "significantly enhanced" * "These results underscore the critical role of [...]" * "Key Aspects", "Key Observations", "Key findings" So why is this bad? - Because it undermines the trust in the the article. We do not know whether the claims are actually true or whether they were just made up by ChatGPT.
- idonotknowwhy 2y ago
  Thanks for articulating something I've noticed happening all over the internet and even in YouTube video scripts. Are claudisms different from gptisms? Why can't these authors tell ChatGPT to write with a different prose and avoid "delve", "crucial", etc? Wizardlm writes these same things in all its answers too

fancyfredbot 2y ago

Can someone knowledgeable please put these numbers in context with a comparison to an H100? For example they say 1053 tokens per second throughput with batch size 4 on llama2 70b. Is that good?

qeternity 2y ago
Where do you see 1053 t/s for L2 70b? From their benchmark table, they achieve 1.8s response time at bs=4 for 256 output tokens, which would be 569 t/s throughput. At bs=1 you are memory bandwidth bound. So as a simple ceiling we divide the memory bandwidth by the model size (70b is roughly 140gb at fp16). That would yield roughly 24 t/s for H100 and 39 t/s for MI300x. I'm not really sure how GEMM optimizations alone would yield these sort of improvements at fp16. With quantization, where dequant overheads start to become compute bound, or at very high batch sizes, it makes sense and there has been a lot of work with various kernels in this area. But I'm a bit confused at fp16 and bs=1. Maybe someone smarter can chime in? EDIT: Ok I watched their video, which shows 39.75 t/s in the unoptimized (which checks out given the above rough maths) and 313.95 t/s for the optimized. I'm not really sure what's going on here.
- fancyfredbot 2y ago
  I found 1053 t/s in cell G18 in the throughput tab of the full results sheet they linked to: https://docs.google.com/spreadsheets/u/0/d/1hHd9riWh-hQJnNOzEI_ZZN1LpdY_bBSbvDOVPdfYslM/htmlview# https://docs.google.com/spreadsheets/u/0/d/1hHd9riWh-hQJnNOz...
  - qeternity 2y ago
    Hmm the post now links to this doc where I got my figures from: https://docs.google.com/spreadsheets/d/1hHd9riWh-hQJnNOzEI_ZZN1LpdY_bBSbvDOVPdfYslM/edit?gid=1853523197#gid=1853523197 https://docs.google.com/spreadsheets/d/1hHd9riWh-hQJnNOzEI_Z... This sheet only lists the response latency, but it does look more recent than the one you listed. In any event, the sheet you link to shows generation speed and throughput (generation speed is cell E18 which shows 527 t/s). G18 is perhaps prefill + gen but even that doesn't add up.

saagarjha 2y ago

I don't usually like to do this, because clearly someone wrote this post (or at the very least they cared enough to have a LLM help draft it :P). Maybe it's just me and my interests not being aligned with the content. But I just found myself kind of disappointed by the content? Like, the post is about how you can do GEMM tuning on AMD GPUs, a subject which is inherently super interesting–there's a lot of nuance to writing optimal kernels, and some of this is expressed in the article, too. Combine that with an architecture that isn't Nvidia? It's an excellent setup for something that would be interesting to read. Which makes the actual conclusion all the more disappointing, IMO. There's nothing about what the actual optimizations are. It's just "oh yeah we tuned our GEMMs and now LLaMA is faster". Like, I get that nobody actually cares about GEMM and they just want tokens to come out of their GPU. But still, that's like writing a blog post about how you can speed up your game with SIMD and then posting some charts of how Cyberpunk 2077 gives you 2x the frame rate now. Ok, but how? I just feel like the interesting part is missing.

xoranth 2y ago
I believe the reason they offer no details about how they tuned the kernels is that the tuning is done by a tool provided by AMD. See here: https://rocm.docs.amd.com/projects/rocBLAS/en/develop/how-to/Programmers_Guide.html#rocblas-gemm-tune https://rocm.docs.amd.com/projects/rocBLAS/en/develop/how-to...
latchkey 2y ago
This article builds upon the work that C&C just released an article on. https://chipsandcheese.com/2024/06/25/testing-amds-giant-mi300x/ https://chipsandcheese.com/2024/06/25/testing-amds-giant-mi3... Elio, the person who did the testing, worked together with C&C, on my system. He then took it to the next level with this gemm work. The docker container is linked on the article now so that anyone can reproduce the results.

deleted 2y ago

[deleted]

Lindon4290 2y ago

Now, I don't have any MI300X, so I can't make any definite claims here. I am hoping someone else can replicate the results shown here or at the least educate me on how this is possible. Good part is the docker container and associated steps are made public - which is pretty cool! Going by the video, the first thing that gave me pause was that a single MI300X is pulling off groq like performance, i.e. 314 tokens/second for a batch size = 1 (bs=1) with a prompt of 256 tokens and generation of 256 tokens. [1] The Llama-2 70B is 128.48GB with FP16 (you can see this in the video). The entire model fits well within the 192GB HBM memory of the MI300X - which is awesome! However, for an regressive transformer model, during generation, the entire model weights are processed to generate a single next token. These models are "next token predictors" so to say, and you need the previous token to generate the next token. Therefore, the 128.48GB of model weights need are consumed from the HBM at the compute cores of the MI300X, per generated token. Note, I am not talking about the prefill - which only needs a single forward pass to generate the first output token. Every subsequent output token is auto-regressive. The video shows that a single prompt (bs=1) with 256 token prompt and 256 tokens generated within 1.63 second. There is no tensor parallelism involved, or batching or anything else. This is a bs=1 case with a single card, and hence you can reason about the math fairly easy. This shouldn't fundamentally be possible within the specs of the MI300X card. The card has a peak HBM memory bandwidth of 5.3 TB/s. You'll notice that to cycle through the weights (assuming FP16) 256 times, you'd need a minimum of 6 seconds, even at perfect ideal conditions. Napkin math: (256 * 128.48e9) / (5.3e12) [1] https://wow.groq.com/groq-sets-new-large-language-model-performance-record-of-300-tokens-per-second-per-user-on-meta-ai-foundational-llm-llama-2-70b/ https://wow.groq.com/groq-sets-new-large-language-model-perf...

latchkey 2y ago
Elio clarified this to me, he said... 256 tokens in 1 prompt/batch Not 256 batches
- Lindon4290 2y ago
  Yes, a single sequence with 256 prompt tokens and 256 output tokens. This is a batch size = 1. No one is saying anything about 256 batches. The first step in understanding this is to notice that the model (llama2) generates 1 output token at a time. This is because the llama2 70B is a autoregressive decoder-only transformer. Fundamentally, to generate a single output token you need to process the entire model weights. At each forward pass you generate 1 token. OK, now to generate 256 output tokens - you need 256 sequential forward passes. At each forward pass, the entire model is read from the gpu VRAM. Even at ideal memory bandwidth (5.3 TB/s) that (256 forward passes of a 128.48GB model) should take 6s. The reported number of 1.63s should not be possible. I'd strongly recommend checking for correctness - that the generate output is coherent. Try sending actual prompts to the "gemm tuned" model and observing the generated responses and latencies. With the "benchmark_throughput.py" you only get a final number and there is no check whether the output is valid or not.

fancyfredbot 2y ago

The numbers in this article don't make sense. They aren't consistent with the hardware (they seem to show the weights being loaded faster than mi300x peak memory bandwidth) and they aren't self consistent (70B models running only 2x slower than 7B models). I'm not sure what's going on here but something is not right.

Lindon4290 2y ago
Yeah, I've asked questions along those lines as well. Something sketchy is going on. See here - https://news.ycombinator.com/item?id=40833109 https://news.ycombinator.com/item?id=40833109
- deleted 2y ago
  [deleted]