I'm not sure what you mean. All CUDA code needs to be aware of the programming model the GPU imposes on them, splitting their code manually into threads and warps and blocks and kernels to match. This isn't really transparent at all.

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.

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, certainly far from the 2x+ it promised.

Ten years on and the AMD 9950X does a pretty good job with it, however.

I mean, 288-E Core Xeons are about to ship. Xeon 6900 series, right? (Estimated to ship in Q1 2025)

So Larrabee lives on for... some reason. These E cores are well known to be modified Intel Atom cores and those were modified Xeon Phi cores which were Larrabee based.

Just with.... AVX512 being disabled. (Lost when Xeon Phi turned into Intel Atoms).

Intels technical strategy is completely bonkers. In a bad way. Intel invented all this tech 10 to 20 years ago but fails to have a cohesive strategy to bring it to market. There's clearly smart people there but somehow all the top level decisions are just awful

Maximum performance? Okay, you'll have to upgrade to ballot instructions or whatever and rearchitect your algorithms. (Or other wavefront / voting / etc. etc. new instructions that have been invented. Especially those 4x4 matrix multiplication AI instructions).

But CUDA -> PTX intermediate code has allowed for significantly more flexibility. For crying out loud, the entire machine code (aka SASS) of NVidia GPUs has been cycled out at least 4 times in the past decade (128-bit bundles, changes to instruction formats, acquire/release semantics, etc etc)

It's amazing what backwards compatibility NVidia has achieved in the past 15 years thanks to this architecture. SASS changes so dramatically from generation to generation but the PTX intermediate code has stayed highly competitive.

NVidia PTX is a very stable interface.

And the PTX to SASS compiler DOES a degree of automatic fine tuning between architectures. Nothing amazing or anything, but it's a minor speed boost that has made PTX just a easier 'assembly-like language' to build on top of.

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.

They said intrinsics. Highway is an abstraction on top of intrinsics.

Not sure why SIMT would help, it requires more compiler transforms than if the code is written for packets/vectors or whatever we want to call them. As you note, cross-lane is a key part of a good SIMD abstraction. Vulkan calls it "subgroups", but from where I sit it's still SIMD.

Square root computation can be tricky, often relying on approximations. These approximations tend to perform best for mid-range values, while accuracy can degrade for very large or very small values. With this in mind, a product of roots is generally more accurate than a root of products.

From a SIMD perspective, it’s worth noting that on most platforms, the cost of computing one square root or two is the same. On modern x86 server CPUs, for instance, you can calculate up to 8 double-precision roots in parallel with identical latency. So there’s no additional cost in terms of performance.

I hope this sheds some light on the design of my code.

PS: In a previous life, I did research in Astro- and Plasma Physics. While I don’t claim to remember all the Math, it’s usually more productive to ask for clarification than to assume ignorance ;)

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

Numerics in .NET are not a high-level abstraction and do out of box what many mature vectorized libraries end up doing themselves - there is significant overlap between NEON, SSE* and, if we overlook vector width, AVX2/512 and WASMs PackedSIMD.

.NET has roughly three vector APIs:

- Vector<T> which is platform-defined width vector that exposes common set of operations

- Vector64/128/256/512<T> which has wider API than the previous one

- Platform intrinsics - basically immintrin.h

Notably, platform intrinsics use respective VectorXXX<T> types which allows to write common parts of the algorithm in a portable way and apply platform intrinsics in specific areas where it makes sense. Also some method have 'Unsafe' and 'Native' variants to allow for vector to exhibit platform-specific behavior like shuffles since in many situations this is still the desired output for the common case.

The .NET's compiler produces competitive with GCC and sometimes Clang codegen for these. It's gotten particularly good at lowering AVX512.

I will respectfully disagree with your statement, with the caveat that I mostly dabble in arithmetic with 128b/256b float and int vectors.

Using C or C++ with vector extensions (Gcc/Clang) or Rust (nightly) std::simd is very easy and you get code that is portable to different CPUs and ISAs.

But most importantly they have a zero cost fallback option to CPU-specific intrinsics when you need them. An f32x8 can be passed at zero cost as __mm256 to any core::arch::x86_64::__mm_intrinsic (or xmmintrin.h in C++ land).

You gain portable arithmetic and swizzles and SIMD vector types, but lose nothing. Not having to write everything for x86_64 and aarch64 is a huge win even if doesn't quite cover everything.

Additionally you can use wider vectors than your hardware supports, the compiler is able to split your f64x64 to 128, 256 or 512 bit registers as needed depending on the compile target.

The library approach does a pretty good job in conjunction with a good compiler, and sensible algorithm design.

You're writing C++ code but as if it was shader code.

I've seen impressive results with clang doing this sort of thing.

There's the middle-ground approach of having primarily target-specific operations but with intersecting ones named the same, and allowing easily building custom abstractions on top of such to paper over the differences how best it makes sense for the given application. That's the approach https://github.com/mlochbaum/Singeli takes.

There's a good amount of stuff that can clearly utilize SIMD without much platform-specificness, but doesn't easily autovectorize - early-exit checks in a loop, packed bit boolean stuff, some data rearranging, probing hashmap checks, some very-short-variable-length-loop things. And while there might often be some parts that do just need to be entirely target-specific, they'll usually be surrounded by stuff that doesn't (the loop, trip count calculation, loads/stores, probably some arithmetic).

I still like ispc, but that's not going to catch on.

I’m not sure about EVE. I trialled it by trying to uppercase a string and even though I got it working in the end it was quite unpleasant. Their docs need to be better.

Unfortunately a bit late :) Highway reached v1.0 about 2.5 years ago. How long would it take until Clang/GCC/MSVC are ready, and all users' distros have updated? Not to mention that the number of ops provided by std::experimental::simd is extremely limited - basically only math operators, and zero support for shuffling/crypto/rounding/interleaving/table lookups which seem indispensable for many applications.

That may have been my mistake. I use super & hyper interchangeably and don't always notice :)

PS: Should be an easy patch, will update!

Patched ;)

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

If you run many small tasks on the GPU, you can increase throughput by overlapping transfers and computation. There may also be other ways to batch problems together, but that depends on the algorithms.

The one truly unfixable issue is round-trip latency.

> The GPU often loses even for small neural nets given the large latency

Apple's neural engine shows that you can live in between those two worlds.

As you said, the trouble is the latency, the programming model is still great.

Apple chips share the same physical memory between the GPU and the CPU. Still, they don't have USM/UVM (Unified Shared Memory/Unified Virtual Memory), that is, the GPU and the CPU can't access the same data concurrently and easily. Programs must map/unmap pages to control which device accesses it, and that's a very expensive operation.

Not much. Synchronisation of tasks is still a big overhead.

If you've got tens to hundreds of microseconds worth of workload, sure, get the GPU to do it.

But bear in mind 1000 clocks at 4GHz is 250ns, there's still a sizeable region where tight CPU/GPU integration isn't tight enough.

I think an argument could be made depending on the real world timings. How much closer in time is the Apple GPU vs one on a PCIe bus?

That's even better then. Just let the AI read the literature and write the optimal compiler.

For humans it's very hard but it will be a breeze for the AI. I thought HN was a community of builders. This is an obvious startup opportunity.