Live data from Hacker News

Understanding SIMD: Infinite complexity of trivial problems

modular.com

91–100 of 127 posts

Re: Understanding SIMD: Infinite complexity of trivial problems

#91

You can simplify the 2x sqrts as sqrt(a*b), overall less operations so perhaps more accurate. It would also let you get rid of the funky lane swivels. As this would only use 1 lane, perhaps if you have multiple of these to normalize, you could vectorize it.

my thoughts exactly - crazy to know all these arcane SIMD opcodes but not know basic maths!!

Moments like that are enlightening. When you see something really improbable (knowing advanced SIMD while appearing to ignore basic algebra), it's likely the moment you see a gap in your picture of the world. So it's tine to learn something new and likely unexpected (else you could have guessed).

Re: Understanding SIMD: Infinite complexity of trivial problems

#93
post #75

Earlier quoted context omitted.

How much of this is because CUDA is designed for GPU execution and because the GPU ISA isn’t a stable interface? E.g. new GPU instructions can be utilized by new CUDA compilers for new hardware because the code wasn’t written to a specific ISA? Also, don’t people fine tune GPU kernels per architecture manually (either by hand or via automated optimizers that test combinations in the configuration space)?

My understanding is that there is a lot of hand-writing (not just fine-tuning) going on. AFAIK CuDNN and TensorRT are written directly as SASS, not CUDA. And the presence of FP8 in H100, but not A100, would likely require a complete rewrite.

Cub, thrust and many other libraries that make those kernels possible don't need to be rewritten.

When you write a merge sort in CUDA, you can keep it across all versions. Maybe the new instructions can improve a few corner cases, but it's not like AVX to AVX512 where you need to rewrite everything.

Ex: https://github.com/NVIDIA/cub/blob/main/cub/device/device_me...

Re: Understanding SIMD: Infinite complexity of trivial problems

#94

Intel needs to see what has happened to their AVX instructions and why NVidia has taken over. If you just wrote your SIMD in CUDA 15 years ago, NVidia compilers would have given you maximum performance across all NVidia GPUs rather than being forced to write and rewrite in SSE vs AVX vs AVX512. GPU SIMD is still SIMD. Just... better at it. I think AMD and Intel GPUs can keep up btw. But software advantage and long te…

While there is some truth in what you say, it makes seem like writing in the CUDA style is something new and revolutionary invented by NVIDIA, which it is not. The CUDA style of writing parallel programs is nothing else than the use of the so-called "parrallel do" a.k.a. "parrallel for" program structure, which has been already discussed in 1963. Notable later evolutions of this concept have been present in "Communic…

Oh all of this was being done in the 1980s by Lisp* programmers.

I'm not calling it new. I'm just saying that the intrinsics style is much much harder than what Lisp*, DirectX HLSL, CUDA, OpenCL (etc. etc) does.

A specialized SIMD language makes writing SIMD easier compared to intrinsic style. Look at any CUDA code today and compare it to the AVX that is in the above article and it becomes readily apparent.

Re: Understanding SIMD: Infinite complexity of trivial problems

#95
post #76

Earlier quoted context omitted.

It's not the same. AVX2 instructions haven't changed and never will change. In contrast, NVidia can go from 64-bit instruction bundles to 128-bit machine code (96-bit instruction + 32-bit control information) between Pascal (aka PTX Compute Capacity 5) and Voltage (aka PTX Compute Capacity 7) and all the old PTX code just autocompiles to the new assembly instruction format and takes advantage of all the new memory ba…

I hope people aren't writing directly to AVX2. When using a wrapper such as Highway, you get exactly this kind of update after a recompile, or even just running your code on a CPU that supports newer instructions. The cost is that the binary carries around both AVX2 and AVX-512 codepaths, but that is not an issue IMO.

> I hope people aren't writing directly to AVX2.

Did you not read the article? It's using AVX intrinsics and NEON intrinsics.

Re: Understanding SIMD: Infinite complexity of trivial problems

#97
post #23

I see a lot of "just use the GPU" and you'd often be right. SIMD on the CPU is most compelling to me due to the latency characteristics. You are nanoseconds away from the control flow. If the GPU needs some updated state regarding the outside world, it takes significantly longer to propagate this information. For most use cases, the GPU will win the trade off. But, there is a reason you don't hear much about systems…

Despite my 'Use a GPU' post below, you are absolutely correct. Maximizing performance on a CPU today requires all the steps in the above article, and the article is actually very well written with regards to the 'mindset' needed to tackle a problem such as this. It's a great article for people aiming to maximize the performance on Intel or AMD systems. ------ CPUs have the memory capacity advantage and will continue…

> CPUs have the memory capacity advantage

perhaps also more precisely they also have quite an advantage on anything that needs and plays nicely with caches? when I sliced my problem to maximize cache usage, I saw pretty clear scalability with cores: L1/L2 cache bandwidth is ~30GB/s, so e.g. a 32 core system starts to compete with the big consumer GPUs.

Re: Understanding SIMD: Infinite complexity of trivial problems

#98
post #70
post #15

Earlier quoted context omitted.

It is worse than that, given that AVX is the survivor from Larrabee great plan to kill GPUs. Larrabee was going to take over it all, as I enjoyed its presentation at GDCE 2009.

And a few years later, Intel said we'd get AVX512 on everything by 2016, and that the instruction encoding supported a future extension to 1024. And then the Skylake and Cannon Lake debacle.. First they pulled it from the consumer chips a fairly short time before launch. Then the server chips it was present in would downclock aggressively when you did use it, so you could get at best maybe 40% more performance, certa…

Oh, and I neglected to mention the protracted development, and short, miserable life, of Cannon Lake itself.

First announced in 2013, it eventually shipped five years later in only a single, crippled dual-core mobile SKU, which lasted just a year in the market before they killed it off.

"Let's put our only consumer implementation of our highest performing vector architecture on a lame-duck NUC chip.", good move guys.

Re: Understanding SIMD: Infinite complexity of trivial problems

#99
post #92

Did they write bfloat16 and bfloat32 when they meant float16 and float32? On the image: https://www.modular.com/blog/understanding-simd-infinite-com...

Yeah I was really confused at first, pretty sure they messed up the labels.

Patched ;)

Re: Understanding SIMD: Infinite complexity of trivial problems

#100

Earlier quoted context omitted.

Despite my 'Use a GPU' post below, you are absolutely correct. Maximizing performance on a CPU today requires all the steps in the above article, and the article is actually very well written with regards to the 'mindset' needed to tackle a problem such as this. It's a great article for people aiming to maximize the performance on Intel or AMD systems. ------ CPUs have the memory capacity advantage and will continue…

> CPUs have the memory capacity advantage perhaps also more precisely they also have quite an advantage on anything that needs and plays nicely with caches? when I sliced my problem to maximize cache usage, I saw pretty clear scalability with cores: L1/L2 cache bandwidth is ~30GB/s, so e.g. a 32 core system starts to compete with the big consumer GPUs.

Caches between CPU and GPU are extremely complex.

Not only because caching is complex in CPU land, but because GPU has a completely different set of caches (and registers) that depends entirely on architecture.

Case in point: GPUs often have access to 256 architectural registers aka 1024-bytes of register space. Now this depends on how much occupancy your GPU code is targeting (maybe 4-occupancy?), but there's a lot you can do with even 64-registers (aka: 4-occupancy and 256-bytes of register space), and is key to making those blazing fast FP16 matrix-multiplication kernels "for AI" everyone's so hot about right now.

For something like FP16 matrix multiply (a very cache-friendly problem), the entire SIMD (all 32-lanes) of the GPU symmetric multiprocessor work together on the problem. So we're talking about an effective 32kB of __register space__ (let alone cache or other memory in the hierarchy).

Even before FP16 matrix-multiply instructions, this absurd register space advantage is why GPUs were the king of matrix multiplication.

--------

GPUs are worse at larger cache sizes, say 1MB or 2MB. At 1MB+, a modern CPU core's L2 cache can hold all of that (either Intel P-core, AMD Zen5, or even Intel E-core can hold many MB in L2).

GPUs have a secret though: a crossbar at the __shared__ memory level. Rearranging memory and data across your lanes can be done through this crossbar (including many-to-one reductions for atomics, or one-to-many broadcasts in just a single clock tick). So your GPU-lanes have incredible communication available to them (and this crossbar is the key for modern ballot / voting based horizontal compute of modern GPU styles). This is only ~64kB of space and shared between all GPU-lanes but with 1024-lanes supporting communication its an important element of GPU memory.

CPU L3 cache is very nice from a latency perspective, but bandwidth wise L3 cache is on the order of GPU's GDDR6x or HBM. 500GB/s to 2000GB/s, depending on the technology.

Finally, CPU DDR5 RAM may be the slowest in this discussion, but its the biggest. 2TB+ Xeon Servers aren't even that expensive and can be assumed for any serious tech firm these days (ie: all-RAM Databases and whatnot).

---------

So at different sizes and different use-scenarios, GPUs and CPUs will trade places. I'd expect CPUs to win most red-and-black tree races, but GPUs will win matrix multiplication. Both take advantage of "cache" but in very different ways.

Post reply on HN