https://twitter.com/abhi9u/status/1715753871564476597

That is against HN's rules. In fact, it's the one thing that's important enough to be in both the site guidelines and FAQ. HN users feel extremely strongly about this.

Q: Can I ask people to upvote my submission?

A: No. Users should vote for a story because they personally find it intellectually interesting, not because someone has content to promote. We penalize or ban submissions, accounts, and sites that break this rule, so please don't.

https://news.ycombinator.com/newsfaq.html

Don't solicit upvotes, comments, or submissions. Users should vote and comment when they run across something they personally find interesting—not for promotion.

https://news.ycombinator.com/newsguidelines.html

But I won't do it again now that I know.

Surprised there's no mention of async copies here. If you want to get the most out of the GPU, you don't want it idle when copying data between the host and the GPU. Many frameworks provide for a mechanism to schedule async copies which can execute along side async work submission.

The post is sort of GPU 101 but there's a whole world of tricks and techniques beyond that once you start doing real-world GPU programming where you want to squeeze as much out of the expensive GPU as possible. Profiling tools help a lot here because, like much of optimizing now, there are hidden cliffs and non-linearities all over that you have to be aware of.

But if you use a GPU with a large number of FP64 units, it may speed things up a lot. These are generally not gaming GPUs, but if you have a 4060 sitting around anyway, it has about 300 GFLOPS FP64 performance, likely more than your CPU. Modern CPUs are mighty in this regard though, able to issue many FP64 operations per clock per core.

maybe this article is brilliant, but when the first line is something so blatantly untrue it really makes it hard to take the rest seriously

Took some philosophy courses for fun in college. I developed a reading skill there that lets me forgive certain statements by improving them instead of dismissing them. My brain now automatically translates over-generalizations and even outright falsehoods into rationally-nearby true statements. As the argument unfolds, those ideas are reconfigured until the entire piece can be evaluated as logically coherent.

The upshot is that any time I read a crappy article, I'm left with a new batch of true and false premises or claims about topics I'm interested in. And thus my mental world expands.

I tried real hard to understand what some continental philosophers, such as Latour, Deleuze, and Žižek are on about, giving some of their texts quite some benefit of the doubt. After about ten years of doing so, I am more and more returning to my previous opinion that some of these just like to talk, and even though they have lots to say, they say so much, that I still have to do all the actual philosophizing myself.

To that end, you have to approach philosophical reading with a degree of hypothetical detachment that not everyone is immediately comfortable with.

Please tell me you forgot the /s.

I have a PhD in computer engineering from a top-20 school in US. Took a bunch of grad level classes, passed the quals (my specialty was ML accelerators).

I do NOT have an “intimate understanding of CPUs”. I probably know a little bit more about CPUs than an average programmer. Which is very little.

Modern CPUs are extremely complex. Almost as much of impenetrable black boxes as modern neural networks.

and 101 other hilarious jokes you can tell yourself!"

In my school to pass the computer architecture course you had to read and present a recent paper on CPU design.

I'm very confident in people having the ability to read by the time they are in college. And considering that summarizing a research paper doesn't even have to be perfect, plus very little need to scrutinize the experiment itself, undergrads should be able to do that.

Otherwise they don't belong in college.

I did an undergrad in CS, where I did well. I don't feel like I understand CPUs very well. Certainly not anywhere in the realm of "intimate."

I think the descriptor “intimate” is overstating the case if you don’t know how to optimize code on different implementations of the same ISA. Most formal CS programs give you a generic understanding of CPUs, more like a survey course, not enough information to do serious optimization.

Compilers are good at finding very local optimizations targeting specific microarchitectures. If you use godbolt and switch out the architecture flags you can see how the codegen changes.

However, as with higher level code, the compiler can’t rewrite your data structures or see patterns spread across too much code. As a simple example, modern cores have 4+ concurrent ALUs with different capabilities. Some of those ALUs are usually idle if there are not enough independent instructions or independent instructions are too far away from the current instruction. You can see large performance improvements simply by reorganizing your C code so that the compiler and CPU can see more opportunities to use more ALUs in parallel. Interestingly, a lot of these code changes are trivial no-ops at the code semantics level, but they bring the ALU concurrency opportunity within view of the CPU.

There is an active niche community on the Internet that studies how various instruction sequences interact with various microarchitectures. This is probably the best resource because a lot of detail is not well documented by the CPU companies themselves. It requires a fair amount of experimentation to develop an intuition for how code will run on a given CPU at this level of detail.

If your day to day involves "intimate knowledge with the CPU" and bumping program counters, I can rule out a lot of things that you probably aren't building, like a compelling iOS app or forum HTTP server, for example.

Maybe you're doing impressive work on emulators or something though. But it's nothing to get pompous about just because other people don't share that interest.

We too easily go off careening into a circlejerk.

https://news.ycombinator.com/item?id=37965142

It could have been worded more constructively sure, but given that the entire point of this website is to have discussions about the material posted, I think it’s added quite a bit.

I learned basic facts about CPU architectures at university, know in a very basic way the landscape of things and occasionaly stumble upon updates to my limited knowledge... but by no means would I say that, rather like "a basic understanding of how CPUs work / are designed / are to be used" (?)

If I were proficient in assembler, maybe I'd claim to have an "intimate understanding" of how to use CPUs at a low level (still sounds a bit braggy)

It still is not the same though as being an expert in CPU/GPU design.

So yeah I agree.

Article is interesting though, esp. the diagram!

However, given the prevalence of the von Neumann computing architecture, I don't think it's completely off the mark - even if people don't know von Neumann's name :)

Wave intrinsics in HLSL, and similar CUDA things can read registers from different threads within the current wavefront.

Also, in the paragraph about memory architecture, I would mention the caches provide no coherency guarantees across threads of the same dispatch/grid, but there’s a special functional block global to the complete chips which implements atomics on global memory.

Want to run a calculation for every pixel on your screen? No problem.

Want to have a branching condition? Ouchie.

Same with drone versus quad-copter, etc...

[0] e.g., https://www.jetsonaero.com/jetson-one

[1] e.g., https://en.wikipedia.org/wiki/HESA_Shahed_136

But I put SIMD in the category of something that isn't necessary once one has learned other (more flexible) paradigms. I prefer MIMD and clusters/transputers, which seem to have died out by the 2000s. Today's status quo puts the onus on developers to move data manually, write shaders under arbitrary limitations on how many memory locations can be accessed simultaneously, duplicate their work with separate languages for GPU/CPU, know if various hardware is available for stuff like ray tracing, get locked into opinionated frameworks like OpenGL/Metal/Vulkan, etc etc etc. GPUs are on a side tangent that can never get me to where I want to go, so my experience over the last 25 years has been of a person living on the wrong timeline. I've commented about it extensively but it just feels like yelling into the void now.

Loosely, a scalable general purpose CPU working within the limitations of the end of Moore's law is multicore with local memories, sharing data through a copy-on-write content-addressable memory or other caching scheme which presents a single unified address space to allow the user to freely explore all methods of computation in a desktop computing setting. It uses standard assembly language but is usually programmed with something higher level like Erlang/Go, Octave/MATLAB or ideally a functional programming language like Julia. 3D rendering and AI libraries are written as a layer above that, they aren't fundamental.

It's interesting that GPUs have arrived at roughly the multicore configuration that I spoke of, but with drivers that separate the user from the bare-metal access needed to do general purpose MIMD. I had thought that FPGAs were the only way to topple GPU dominance, but maybe there is an opportunity here to write a driver that presents GPU hardware as MIMD with a unified memory. I don't know how well GPU cores handle integer math, but that could be approximated with the 32 bit int portion of a 64 bit float. Those sorts of tradeoffs may result in a MIMD machine running 10-100 times slower than a GPU, but still 10-100 times faster than a CPU. But scalable without the over-reliance on large caches and fast busses which stagnated CPUs since around 2007 when affordability and power efficiency took priority over performance due to the mobile market taking over. And MIMD machines can be clustered and form distributed compute networks like SETI@home with no changes to the code. To get a sense of how empowering that could be to the average user: it's like comparing BitTorrent to FTP, but for compute instead of data.

Looking at this quote:

> the Nvidia H100 GPU has 132 SMs with 64 cores per SM, totalling a whopping 8448 cores.

8448 cores sure sounds impressive. But the Apple M2 Ultra only has 76 cores?!

How can the NVidia H100 GPU have over 110x more cores? Clearly it doesn’t have 110x more performance over the M2 Ultra, so what is going on here?

Yes, their architecture is different from many in that they support a separate program counter per lane (which is why they feel justified in calling this a “thread”), but ultimately it’s the rate and throughput of ALUs that matter.

I ask because I’ve seen posts from NVIDIA support saying that divergence is still very expensive and I’ve also seen benchmarks that force divergence in each warp by evenly splitting the warp, and the benchmarks result in 2x runtime when that happens vs. when the control-flow is dynamically uniform.

One thing to keep in mind is that even if you were to dynamically reform-warps, there’s still a potential expense because you then lose the advantage of doing things like accessing adjacent elements of data in adjacent threads. You’re bound to now have more bank conflicts, fewer memory accesses being coalesced, etc. Perhaps they do actually do this warp re-formation, but that itself does have this additional cost.

See this diagram from an nvidia blog post: https://developer-blogs.nvidia.com/wp-content/uploads/2021/g...

( https://developer.nvidia.com/blog/nvidia-ampere-architecture... )

Another piece in the "why is ML so inefficient" puzzle!

I wonder what that memory copying overhead is IRL. If it's like normal stuff it'll be brutal. I mean, they offload tcp processing into hardware to avoid that. This is way more data, though it is done in bigger chunks.

But really, I wondered why the cpu-bound LLMs do quantization, which from what I understand is the process of reducing the precision of the weights to use less memory.

Does the lack of precision make a difference? It's unclear. If that's the case, then why use FP at all? If precision doesn't matter, then the extra precision is just making the process use more resources for no real reason. And likely orders of magnitude more resources than required.

I mean, this field wasn't started by people who understood performance. They used tools and built something...but there's no 'why.' This is what the tools did, so they did that.

Here's why that might be important: on a normal CPU, accessing data one way can be orders of magnitude faster than another way...but you have to be aware. Would you like to reduce your LLM costs by orders of magnitude?

For shallower networks, there can be an advantage to copying just the original compressed data to GPU memory and then doing decompression etc there. But modern GPUs haven’t adopted pcie 5 yet because the raw compute is more important.

Lastly tensor cores have had a big impact, depending on the network they can be fast enough to be very under-utilized.

I'm not trying to be snarky. I think there's an unhelpful compunction to want to know everything about everything among STEM types like programmers (of which I am one). Specialization is fundamental to the success, not just of whole economies, but of the individuals in them. It can feel like a paintful sacrifice to admit that you'll never (have time to) learn, say, the entire Python language specification, or how type inference works, or any number of other things someone might tell you is critical knowledge. But it's often liberating, and more often than that, mandatory.

(Maybe I was trying to be a little snarky initially, but I'm not any more.)

And even if you do have time, you'll probably forget most of it if you don't make use of it in your daily job. Also if you do want to learn about a new topic, it'll take commitment. I remember going through an MIT OS project, it took me something like 50 hours of hours to complete the project. Pretty much impossible when you have a full time job (I didn't at the time). And despite that, I still consider myself a newbie in OS development.

That being said, a little extra knowledge can come handy and make a difference in an interview for instance, or reduce ramp-up time when changing teams.

This is also what school is for. Give you full time and a structured program to pick the fundamentals. There's only so much you can learn once you have a demanding job. It's actually pretty sad we don't get to go back to school in the middle of our career.

Edit: still an interesting article ;)

I took an OS course loosely based on MIT one, and of course wouldn't consider myself an experienced OS developer as well, but I think the residual knowledge is sometimes really helpful when debugging bad application performance at work.

Things every developer should know about English, part 1: why word order matters in negation /s

Obviously you meant "not every developer needs to know" instead of "every developer doesn't need to know", but I see this switch so often lately that I'm beginning to wonder if it's a dialect thing (similar to double negative implying a positive to some and an emphasis of the negative to others)

> [Foreach developer D, D] needs to know this.

> [Foreach developer D, D] doesn't need to know this. (No developers.)

But you could also have:

> [The logical conjuction of all developers[0]] needs to know this.

> [The logical conjuction of all developers] doesn't need to know this. (Not all developers.)

The former convention is clearly better unless I've missed something rather significant, but inherently broken features in a number of programming languages (eg perl 6 'junctive' operators) suggest the latter is well-established as a thing that exists.

0: ie, a linguistic fiction that has only those properties / needs only those things that are common to every developer.

But in this case it doesn't even mean that: all of this applies to a substantial minority of programmers, but definitely a minority.

The first being number base systems.

Yes, you want to know as much as possible ( helps debugging/zooming in on issues since you don’t need to introduce an outsider to your problem, helps avoid errors), yes you need to specialize somewhere, no you can’t know everything and often don’t need to

Even general knowledge has depth. Even if you are generalist, or specialist in different area you should deepen your knowledge at every area gradually.

At big companies you have SWE L1, L2, L3, senior, staff, principal, etc. You also have SRE and maybe some devops or architect roles.

Smaller companies you have lots of people generic “engineer” titles or “full stack engineers”, etc.

Why don’t we encode people’s specialties in their titles if most engineers are working on narrow sections of software?

E.g. “React & Node Developer” instead of “Full stack engineer” … etc

I suppose the easiest rationale is generic titles allow for easier mobility between disciplines.

Edit: I wanted to mention that IME in STEM, most? More often than not? are goal oriented… It is _completely_ okay to not have any “working” thing at the end of your tinkering/learning… The journey into that “thing” can be a learning experience in and of itself. I’ve often started to learn things but after a certain point have told myself “I’m gonna stop here and that’s okay. I don’t need to have a solid understanding of this, at this time”. YMMV…

Now, about every developer. Ideally, developers should know something, in general, about all fields related to programming. Similar to how in medicine you need to learn about different branches of medicine, even though you'll specialize in one, and so it is in mathematics, physics and so on.

Presently, the demand for quality professionals in programming is very low. There aren't any good testing or certification programs that can tell a good programmer from a bad one. The industry is generally happy with "specialists" who perhaps only know to do one thing, somewhat. So, presently you don't need to know anything about GPU or any other field that's not directly related to your job description.

----

Now, about the article itself. While it gives a lot of valuable factual information, it's missing the forest for the trees. It's very dedicated to how CUDA works or some other particular aspects of NVidia's GPUs. The part that's missing is the part that could make it, potentially, a candidate for the kind of introduction to GPU programming that would make it worth reading to expand your general understanding of how computers work.

If you ever paid attention to how encyclopedic articles are written: the structure of a definition given by encyclopedia often has two components. First puts the object of the definition into the more general category, second explains how the object of the definition is different from other elements of the category. What the GPU article is missing is the first component: putting GPU programming into the more general category. This, in practical terms, means that questions like "is DPU programming anything like GPU programming?" or "can `smart' SSDs (with FPGAs) be treated similar to GPUs?" unanswered.

The main reason GPU gets so much attention these days is that CPU manufacturers (Intel in particular) simply cannot get their act together. Intel had promised significant breakthroughs with Xeon Phi like a decade ago.

In the meantime people have invented more and more applications that need significant computational power. But it will eventually get there. E.g., AMD' latest epyc features 96 cores. Importantly, that computational power is available in principle with simpler / more familiar programming models.

Then I use a vectorized haversine algorithm on these arrays to obtain a third one with the distances between each "row" of the two arrays.

With NumPy this is fast, but I guess that a GPU could also perform this and likely do it much faster. Would it be worth it if one has to consider that the data needs to be moved to the GPU and the result be retrieved?

This is implemented in an nvidia geospatial library call cuspatial: https://docs.rapids.ai/api/cuspatial/legacy/api_docs/spatial...

> I duplicate the array and remove the first item in the duplicated array, append the last entry of the original array to the duplicate in order for them to be equal in length.

I assume/hope this is only what you're doing logically, not physically?

Otherwise you might as well just compute the length using n - 1 points from the existing array, then do the remaining portion manually and add it to the existing sum. That would avoid the copying of the whole array.

Depends on your batch size. If the computation on the CPU is less than lets say 200 ms it's probably not worth it.

Also consider that integrated GPUs don't have separate memory, I'm not sure, but they might not have a high cost of moving data to memory.

I've already refrained myself to learn vulkan because it scares me, but similarly, opengl and cuda are a bit mysterious to me, and I don't really know how I could take advantage of it, since most computing tasks cannot be made parallel.

I've read there are data structures that are somehow able to take advantage of a GPU as an alternative to the CPU (for example a database running on a GPU), but it seems a very new domain, and I don't have the skill to explore it.

Writing and launching raw cuda kernels is too low level for me, but writing with Thrust makes it feel pretty similar to writing regular C++ code. You still need to deal with moving data from host to device and back, but that's as simple as assigning a `thrust::device_vector` to a `thrust::host_vector`

I haven’t formally studied any queuing theory, but I think:

1. The rule assumes you have enough processing power to service the average load (otherwise it fails catastrophically like you mentioned)

2. The rule is trying to model the fluctuations in the pending load (which might determine wait time or whatever else).

I'd say the mental model for most programmers is: lines of text, in their language of choice, zipping by really fast.

every developer should know about more than nvidia's spin.

Alexander Titov — Know your hardware: CPU memory hierarchy https://youtu.be/QOJ2hsop6hM

https://github.com/alexander-titov/public/blob/master/confer...

Know Your Hardware - CPU Memory Hierarchy -- Alexander Titov -- C%2B%2B Moscow Meetup March 2019.pdf

https://github.com/alexander-titov/public/blob/master/confer...

GPGPU - what it is and why you should care -- Alexander Titov -- CoreHard 2019.pdf

The main bottleneck isn't compute, it is memory. If you go to talks you're gonna see lots of figures like this one[0] (typically also showing disk speeds, which are crazy small).

Compute is increasing so fast that at this point we finish our operations long faster than it takes to save those simulations or even create the visualizations and put on disk. There's a lot of research going into this, with a lot of things like in situ computing (asynchronous operations, often pushing to a different machine, but needing many things like flash buffers. See ADIOS[1] as an example software).

What I'm getting at here is that we're at a point where we have to think about that IO bottleneck, even for non-high performance systems. I work in ML now, which we typically think of as compute bound, but being in the generative space there are still many things where the IO bottlenecks. This can be loading batches into memory, writing results to disk, or communication between distributed processes. It's one beg reason we typically want to maximize memory usage (large batches).

There's a lot of low hanging fruit in these areas that aren't going to be generally publishable works but are going to have lots of high impact. Just look at things like LLaMA CPP[2], where in the process they've really decreased the compute time and memory load. There's also projects like TinyLLaMa[3] who are exploring training a 1B model and doing so on limited compute, and are getting pretty good results. But I'll tell you from personal experience, small models and limited compute experience doesn't make for good papers (my most cited work did this and has never been published, gotten many rejections for not competing with models 100x it's size, but is also quite popular in the general scientific community who work with limited compute). Wfiw, companies that are working on applications do value these things, but it is also noise in the community that's hard to parse. Idk how we can do better as a community to not get trapped in these hype cycles, because real engineering has a lot of these aspects too, and they should be (but aren't) really good areas for academics to be working in. Scale isn't everything in research, and there's a lot of different problems out there that are extremely important but many are blind to.

And one final comment, there's lots of code that is used over and over that are not remotely optimized and can be >100x faster. Just gotta slow down and write good code. The move fast and break things method is great for getting moving but the debt compounds. It's just debt is less visible, but there's so much money being wasted from writing bad code (and LLMs are only going to amplify this. They were trained on bad code after all). And no, pytorch isn't going to optimize everything for you automagically (no one can. Optimization is often situationally dependent). You're still gonna need to understand the distributed package and it is worth learning MPI.

[0] https://drivenets.com/wp-content/uploads/2023/05/blog-networ...

[1] https://github.com/ornladios/ADIOS2

[2] https://github.com/ggerganov/llama.cpp

[3] https://github.com/jzhang38/TinyLlama

OpenCL was almost unusable due to how difficult it was to set up, plus limited support. Vulcan is terribly ugly and verbose and provides no benefits for pure compute workloads. Writing a "hello word" is a few thousands lines of code, debugging impossible.

Working with CUDA is pure pleasure - immediate setup, super easy CPU/GPU interop and code sharing, all the modern (at the time!) C++ features, super pleasant debugging with stepping, profilers, everything. :)

I totally wish others invested as much into tools as NVIDIA did - but they didn't and made the whole experience miserable. :( If your velocity is 1/10 because of terrible tooling, you need to expect some insane benefits to pick it.

AMD's primary focus for a GPU software ecosystem these days seems to be implementing CUDA with s/cuda/hip

Maybe it’s just where I work, but I feel like even this isn’t true. A lot of the newer/young employees don’t even seem to have a lot of OS/system understanding. A lot of the higher-level languages abstract the immediate need of having any “intimate understanding” of CPUs.

I had done work in the C/C++ code modifying Chromium directly and went, “that’s not how the browser talks to the OS to use the GPU for acceleration”. They immediately tried to argue I was wrong and made some more nonsense up.

This person is a very competent coder, but we got into further discussions and it was made very clear they had lived in a JavaScript landscape for either far too long, or had only barely gone over basic computing concepts. Their understanding of computers largely ended at “and then v8 does some stuff”.

It was weird cause this person had an honours level computer science degree from a very prestigious university and they were not from some rich family that could have paid their way through schooling.

They also had over a decade of experience shipping numerous projects. It really threw me for a loop as to what a software engineer needs to know to succeed at very high levels professionally.

Huh? If they'd got private tutoring, that wouldn't make them understand things any less?

Sure--until he isn't. And that happens north of "senior".

I've had a lot of I-shaped people reporting to and taking direction from me in my career, and I'm not as deep at any one thing as many of them. But I'm deeper at a lot of things than most of those folks, and I can synthesize solutions to more complex problems because of them. Heck, things I learned writing games in my teens still come up on occasion twenty years later.

https://news.ycombinator.com/item?id=37965142

On the other hand I agree that people who take the non-CS route may not be that intimately familiar with all the details.

This is largely irrelevant to the concept of sequential programming.

Not at all. Very few programmers have an intimate understanding of CPUs. I'm not one of them, for example. And that's without mentioning how most programmers actually write Excel scripts, and most of the rest write interpreted code like Javascript or Python which is somewhat removed from the CPU.

> because they grow up writing code for the CPU

There is not really such thing as "The CPU". The design space of CPUs is large enough for it to contain quite a variation of beasts. For example,

> CPUs were designed to execute sequential instructions

Originally, yes, but - one cannot really make this argument about CPUs since Dennard scaling hit the wall, and scale-out began in earnest.

-----

While the mental model is indeed different, hardware-wise, you should think of GPUs as CPUs, except that...

* You only ever use SIMD registers: 32 lanes x 32bit

* Large register file (up to 256 SIMD registers)

* No branch prediction

* Predicated execution supported, but no if-then-else (sort of).

* Very little stack to speak of

* ~10x L1 cache latency

* A piece of L1 "cache" which doesn't cache anything and you can just write to it like regular memory.

* ~100x L2 cache latency, no coherence, and it's smaller

* ~100x main memory latency (and no L3)

* A bunch of special registers for implementing geometric thread & lane indexing.

* ~10x memory bandwidth

* ~5x-10x cores than on a CPU

(numbers are for NVIDIA cards, AMD are slightly different; and of course individual cards differ etc.)

-----

Now take all of those hardware differences, and conceive a programming model. What you get is the fanciful tale of a "not-like-a-CPU" processor.

Programming for the CPU, we often do the opposite, i.e. ignore the cores we have, ignore things like latency hiding, careful management of what's in registers and what's in the cache, and just write naive serial programs.