Don't "optimize" conditional moves in shaders with mix()+step() | Hacker News Reader

Don't "optimize" conditional moves in shaders with mix()+step()

iquilezles.org

409 pointsromes231 comments

I'm sure TFA's conclusion is right; but its argument would be strengthened by providing the codegen for both versions, instead of just the better version. Quote:

"The second wrong thing with the supposedly optimizer [sic] version is that it actually runs much slower than the original version [...] wasting two multiplications and one or two additions. [...] But don't take my word for it, let's look at the generated machine code for the relevant part of the shader"

—then proceeds to show only one codegen: the one containing no multiplications or additions. That proves the good version is fine; it doesn't yet prove the bad version is worse.

The main point is that the conditional didn't actually introduce a branch.

Showing the other generated version would only show that it's longer. It is not expected to have a branch either. So I don't think it would have added much value

Unless you’re writing an essay on why you’re right…

> Unless you’re writing an essay on why you’re right…

He's writing an essay on why they are wrong.

"But here's the problem - when seeing code like this, somebody somewhere will invariably propose the following "optimization", which replaces what they believe (erroneously) are "conditional branches" by arithmetical operations."

Hence his branchless codegen samples are sufficient.

Further, regarding.the side-issue "The second wrong thing with the supposedly optimizer [sic] version is that it actually runs much slower", no amount of codegen is going to show lower /speed/.

The other either optimizes the same, or has an additional multiplication, and it's definitely less readable.

Correct: it would show proof instead of leaving it up to the reader to believe them.

But it's possible that the compiler is smart enough to optimize the step() version down to the same code as the conditional version. If true, that still wouldn't justify using step(), but it would mean that the step() version isn't "wasting two multiplications and one or two additions" as the post says.

(I don't know enough about GPU compilers to say whether they implement such an optimization, but if step() abuse is as popular as the post says, then they probably should.)

Okay but how does this help the reader? If the worse code happens to optimize to the same thing it's still awful and you get no benefits. It's likely not to optimize down unless you have fast-math enabled because the extra float ops have to be preserved to be IEEE754 compliant

..how is it awful if it has the same result?

Because it perpetuates a misconception and is harder to read.

Just look at it.

Fragment and vertex shaders generally don't target strict IEEE754 compliance by default. Transforming a * (b ? 1.0 : 0.0) into b ? a : 0.0 is absolutely something you can expect a shader compiler to do - that only requires assuming a is not NaN.

You missed the second part where article says that "it actually runs much slower than the original version", "wasting two multiplications and one or two additions", based on idea that compiler is unable to do a very basic optimization, implying that compiler compiler will actually multiply by one. No benchmarks, no checking assembly, just straightforward misinformation.

There are 10 types of people in this work. Those who can extrapolate from missing data, and

Making assumptions about performance when you can measure is generally not a good idea.

and what? AND WHAT?

Generated code for RDNA 1:

https://shader-playground.timjones.io/5d3ece620f45091678dcee...

I wish there was a good way of knowing when an if forces an actual branch rather than when it doesn't. The reason people do potentially more expensive mix/lerps is because while it might cost a tiny overhead, they are scared of making it a branch.

I do like that the most obvious v = x > y ? a : b; actually works, but it's also concerning that we have syntax where an if is some times a branch and some times not. In a context where you really can't branch, you'd almost like branch-if and non-branching-if to be different keywords. The non-branching one would fail compilation if the compiler couldn't do it without branching. The branching one would warn if it could be done with branching.

> it's also concerning that we have syntax where an if is some times a branch and some times not.

That's true on scalar CPUs too though. The CMOV instruction arrived with the P6 core in 1995, for example. Branches are expensive everywhere, even in scalar architectures, and compilers do their best to figure out when they should use an alternative strategy. And sometimes get it wrong, but not very often.

For scalar CPUs, historically CMOV used to be relatively slow on x86, and notably for reliable branching patterns (>75% reliable) branches could be a lot faster.

cmov also has dependencies on all three inputs, so if there's a high level of bias towards the unlikely input having a much higher latency than the likely one a cmov can cost a fair amount of waiting.

Finally cmov were absolutely terrible on P4 (10-ish cycles), and it's likely that a lot of their lore dates back to that.

The good way is to inspect the code :)

> it's also concerning that we have syntax where an if is some times a branch and some times not.

It would be more concerning if we didn't. We might get a branch on one GPU and none on another.

>> The good way is to inspect the code :)

The best way is to profile the code. Time is what we are after, so measure that.

For sure.

I think that capability in the shader language would be interesting to have. One might even want it to two-color all functions in the code. Anything annotated nonbranching must have if statements compile down to conditional moves and must only call nonbranching functions.

This is also very relevant for cryptography use cases, where branching is a potential side channel for leaking secret information.

godbolt has rga compiler now, you can always paste in hlsl and look at the actual rdna instructions that are generated (what GPU actually runs, not spirv)

Exactly. Consider this example:

  a = f(z);
  b = g(z);
  v = x > y ? a : b;

Assuming computing the two function calls f() and g() is relativelly expensive, it becomes a trade-off whether to emit conditional code or to compute both followed by a select. So it's not a simple choice, and the decision is made by the compiler.

This is a GPU focused article.

The GPU will almost always execute f and g due to GPU differences vs CPU.

You can avoid the f vs g if you can ensure a scalar Boolean / if statement that is consistent across the warp. So it's not 'always' but requires incredibly specific coding patterns to 'force' the optimizer + GPU compiler into making the branch.

It depends. If the code flow is uniform for the warp, only side of the branch needs to be evaluated. But you could still end up with pessimistic register allocation because the compiler can’t know it is uniform. It’s sometimes weirdly hard to reason about how exactly code will end up executing on the GPU.

f or g may have side effects too. Like writing to memory. Now a conditional has a different meaning.

You could also have some fun stuff, where f and g return a boolean, because thanks to short circuit evaluation && || are actually also conditionals in disguise.

Side effects will be masked, the GPU is still executing exactly the same code for the entire workgroup.

>The reason people do potentially more expensive mix/lerps is because while it might cost a tiny overhead, they are scared of making it a branch.

And the reason for that is the confusing documentation from NVidia and its cg/CUDA compilers. I believe they did not want to scare programmers at first and hid the execution model, talking about "threads" and then they kept using that abstraction to hype up their GPUs ("it has 100500 CUDA threads!"). The result is people coding for GPUs with some bizarre superstitions though.

You actually want branches in the the code. Those are quick. The problem is that you cannot have a branch off a SIMD way so, instead of a branch the compiler will emit code for both branches and the results will be masked out based on the branch's condition.

So, to answer your question - any computation based on shader inputs (vertices, computer shader indices and what not) cannot and won't branch. It will all be executed sequentially with masking. Even in the TFA example, both values of ? operator are computed, the same happens with any conditional on an SIMD value. There can be shortcut branches emitted by the compiler to quickly bypass computations when all ways are the same value but in general case everything will be computed for every condition being true as well as being false.

Only conditionals based on scalar registers (shader constants/unform values) will generate branches and those are super quick.

You can always have the compiler dump the assembly output so you can examine it. I suspect few do that.

Does this also apply for shaders? And is it even useful given the enormous variation in hardware capabilities out there. My impression was that it's all JIT compiled unless you know which hardware you're targeting, e.g. Valve precompiling highly optimized shaders for the Steam Deck

(I'm not a grapics programmer, mind you, so please correct any misunderstandings on my end)

It's all JIT'd based on the specific driver/GPU, but the intermediate assembly language is sufficient to inspect things like branches and loop unrolling.

Not really. DXIL in particular will still have branches and not care much about unrolling. You need to look at the assembly that is generated. And yes, that depends on the target hardware and compiler/driver.

You will have to check for the different GPUs you are targetting. But GPU vendors don't start from scratch for each hardware generation so you will often see similar results.

> So, to answer your question - any computation based on shader inputs (vertices, computer shader indices and what not) cannot and won't branch.

It can do an actual branch if the condition ends up the same for the entire workgroup - or to be even more pedantic, for the part of the workgroup that is still alive.

You can also check that explicitly to e.g. take a faster special case branch if possible for the entire workgroup and otherwise a slower general case branch but also for the entire workgroup instead of doing both and then selecting.

And this is why I wrote There can be shortcut branches emitted by the compiler to quickly bypass computations when all ways are the same value but in general case everything will be computed for every condition being true as well as being false.

I'll comment this here as I got downvoted when I made the point in a standalone comment - this is mostly an academic issue, since you don't want to use step of pixel-level if statements in your shader code, as it will lead to ugly aliasing artifacts as the pixel color transitions from a to b.

What you want is to use smoothstep which blends a bit between these two values and for that you need to compute both paths anyway.

It's absurd to claim that you'd never use step(), even in pixel shaders (there are all kinds of shaders not related to pixels at all).

>since you don't want to use step of pixel-level if statements in your shader code

The observation relates to pixel shaders, and even within that, it relates to values that vary based on pixel-level data. In these cases having if statements without any sort of interpolation introduces aliasing, which tends to look very noticeable.

Now you might be fine with that, or have some way of masking it, so it might be fine in your use case, but most in the most common, naive case the issue does show up.

I don't know how many graphics products you shipped and when, but, say, clamping values at 0, is pretty common even in most basic shaders. It's not magic and won't introduce "aliasing" just for the fact of using it. On the other hand, for example, using negative dot products in you lighting computation will introduce bizarre artifacts. And yes, everyone uses various forms of MSAA for the past 15 years or so even in games. Welcome to the 21st century.

The way you write seems to imply you have professional experience in the matter, which makes it very strange you're not getting what I'm writing about.

Nobody ever talked about clamping - and it's not even relavant to the discussion as it doesn't introduce discontinuity that can cause aliasing.

What I'm referring to is shader aliasing, which MSAA does nothing about - MSAA is for geometry aliasing.

To illustrate what I'm talking about with, an example that draws a red circle on a quad:

The bad version:

    gl_FragColor = vec4(vec3(1.0 - step(0.25, distance(vUv, vec2(0.5)))) * vec3(1.0, 0.0, 0.0), 1.0);

The good version:

    gl_FragColor = vec4(vec3(1.0 - smoothstep(0.24, 0.25, distance(vUv, vec2(0.5)))) * vec3(1.0, 0.0, 0.0), 1.0);

The first version has a hard boundary for the circle which has an ugly aliased and pixelated contour, while the latter version smooths it. This example might not be egregious, but this can and does show up in some circumstances.

Execution with masking is pretty much how broaching works on GPUs. What’s more relevant however is that conditional statements add overhead on terms of additional instructions and execution state management. Eliminating small branches using conditional moves or manual masking can be a performance win.

No, branching works on GPU just like everywhere else - the instruction pointer gets changed to another value. But you cannot branch on a vector value unless every element of the vector is the same, this is why branching on vector values is a bad idea. However, if your vectorized computation is naturally divergent then there is no way around it, conditional moves are not going to help as they also will evaluate both branches in a conditional. The best you can do is to arrange it in such a way that you only add computation instead of alternating it, i.e. you do if() ... instead of if() ... else ... then you only take as long as the longest path.

This reminds me that people who believe that GPU is not capable of branches do stupid things like writing multiple shaders instead of branching off a shader constant e.g. you have some special mode, say x-ray vision, in a game and instead of doing a branch in your materials, you write an alternative version of every shader.

One can do the precisely how it's done in the article -- inspect the assembly.

But you don't generally need to care if the shader code contains a few branches, modern GPUs handles those reasonably well, and the compiler will probably make a reasonable guess about what is fastest.

You do need to care about large non-uniform branches as in the general case the GPU will have to execute both sides.

A non-branching version of the same algorithm will also run code equivalent to both branches. The branching version may sometimes skip one of the branches, the non-branching version can't. So if the functionality you want is best described by a branch, then use a branch.

You got this the wrong way around: For GPUs conditional moves are the default and real branches are a performance optimization possible only if the branch is uniform (=same side taken for the entire workgroup).

The good way of knowing is to look at the assembly generated by the compiler. Maybe not a completely satisfying answer given that the result is heavily vendor-dependent, but unless the high-level language exposes some way of explicitly controlling it, then assembly is what you got.

> The reason people do potentially more expensive mix/lerps is because while it might cost a tiny overhead, they are scared of making it a branch.

This is a problem, though. People shouldn't do things potentially, they should look at the actual code that is generated and executed.

A lot of the myth that "branches are slow on GPUs" is because, way back on the PlayStation 3, they were quite slow. NVIDIA's RSX GPU was on the PS3; it was documented that it was six cycles IIRC, but it always measured slower than that to me. That was for even a completely coherent branch, where all threads in the warp took the same path. Incoherent branches were slower because the IFEH instruction took six cycles, and the GPU would have to execute both sides of the branch. I believe that was the origin of the "branches are slow on GPUs" myth that continues to this day. Nowadays GPU branching is quite cheap especially coherent branches.

coherent branches are "free" but the extra instructions increase register pressure. that's the main reason why dynamic branches are avoided, not that they are inherently "slow".

If someone says branching without qualification, I have to assume it’s incoherent. The branching mechanics might have lower overhead today, but the basic physics of the situation is that throughput on each side of the branch is reduced to the percentage of active threads. If both sides of a branch are taken, and both sides are the same instruction length, the average perf over both sides is at least cut in half. This is why the belief that branches are slow on GPUs is both persistent and true. And this is why it’s worth trying harder to reformulate the problem without branching, if possible.

These sort of avoid-branches optimizations were effective once upon a time as I profiled them on the XBox 360 and some ancient Intel iGPUs, but yeah - don't do this anymore.

Same story for bit extraction and other integer ops - we used to emulate them with float math because it was faster, but now every GPU has fast integer ops.

> now every GPU has fast integer ops.

Is that true and to what extent? Looking at the ISA for RDNA2[0] for instance - which is the architecture of both PS5 and Xbox Series S|X - all I can find is 32-bit scalar instructions for integers.

[0] https://www.amd.com/content/dam/amd/en/documents/radeon-tech...

You're likely going to have a rough time with 64-bit arithmetic in any GPU. (At least on Nvidia GPUs, the instruction set doesn't give you anything but a 32-bit add-with-carry to help.) But my understanding is that a lot of the arithmetic hardware used for 53-bit double-precision ops can also be used for 32-bit integer ops, which hasn't always been the case.

I'm less concerned about it being 32-bit and more about them being exclusively scalar instructions, no vector instructions. Meaning only useful for uniforms, not thread-specific data.

[Update: I remembered and double checked. While there are only scalar 32-bit integer instructions you can use 24-bit integer vector instructions. Essentially ignoring the exponent part of the floats.]

The programming model is that all threads in the warp / thread block run the same instruction (barring masking for branch divergence). Having SIMD instructions at the thread level is a rarity given that the way SIMD is implemented is across warps / thread blocks (groups of warps). It does exist, but only within 32-bit words and really only for limited use cases, since the proper way to do SIMD on the GPU is by having all of the threads execute the same instruction:

https://docs.nvidia.com/cuda/parallel-thread-execution/index...

Note that I am using the Nvidia PTX documentation here. I have barely looked at the AMD RDNA documentation, so I cannot cite it without doing a bunch of reading.

I know all of that. I was talking about RDNA2, which is AMD. There, instructions come in two flavours:

1. Scalar - run once per thread group, only acting on shared memory. So these won't be SIMD.

2. Vector - run across all threads, each threads accesses its own copy of the variables. This is what you typically think of GPU instructions doing.

That does sound like it would be a pretty big limitation. But there appear to be plenty of vector instructions for 32-bit integers in RDNA2 and RDNA3 [0] [1]. They're named V_*_U32 or V_*_I32 (e.g., V_ADD3_U32), even including things like a widening multiply V_MAD_U64_U32. The only thing missing is integer division, which is apparently emulated using floating-point instructions.

[0] https://www.amd.com/content/dam/amd/en/documents/radeon-tech..., p. 259, Table 83, "VOP3A Opcodes"

[1] https://www.amd.com/content/dam/amd/en/documents/radeon-tech..., p. 160, Table 85, "VOP3 Opcodes"

The PTX ISA for Nvidia GPUs supports 64-bit integer arithmetic:

https://docs.nvidia.com/cuda/parallel-thread-execution/index...

It needs to support 64-bit integer arithmetic for handling 64-bit address calculations efficiently. The SASS ISA since Volta has explicit 32I suffixed integer instructions alongside the regular integer instructions, so I would expect the regular instructions to be 64-bit, although the documentation leave something to be desired:

https://docs.nvidia.com/cuda/cuda-binary-utilities/index.htm...

Hmm, it looks like the underlying SASS does have 64-bit integer instructions now, but only with the 12.0 capability level in the recent Blackwell processors. Older versions emulate it via chained 32-bit instructions. Take this example kernel:

  __global__ void add(uint64_t *res, uint64_t x) {
      *res = x + 0x12345;
  }

Compiled with -arch=sm_120, I get the SASS:

  /*0000*/  LDC R1, c[0x0][0x37c]                      ?trans8;
  /*0010*/  LDC.64 R2, c[0x0][0x388]          &wr=0x0  ?trans1;
  /*0020*/  LDCU.64 UR4, c[0x0][0x358]        &wr=0x1  ?trans7;
  /*0030*/  LDC.64 R4, c[0x0][0x380]          &wr=0x1  ?trans1;
  /*0040*/  IADD.64 R2, R2, 0x12345           &req={0} ?WAIT6_END_GROUP;
  /*0050*/  STG.E.64 desc[UR4][R4.64], R2     &req={1} ?trans1;
  /*0060*/  EXIT                                       ?trans5;
  /*0070*/  BRA 0x70;

But with -arch=sm_100, the IADD.64 is broken up into a UIADD3 and a UIADD3.X (contrast with the IADD3 that a regular 32-bit addition would produce):

  /*0000*/  LDC R1, c[0x0][0x37c] ;
  /*0010*/  LDCU.64 UR4, c[0x0][0x388] ;
  /*0020*/  LDC.64 R2, c[0x0][0x380] ;
  /*0030*/  LDCU.64 UR6, c[0x0][0x358] ;
  /*0040*/  UIADD3 UR4, UP0, UPT, UR4, 0x12345, URZ ;
  /*0050*/  UIADD3.X UR5, UPT, UPT, URZ, UR5, URZ, UP0, !UPT ;
  /*0060*/  MOV R4, UR4 ;
  /*0070*/  MOV R5, UR5 ;
  /*0080*/  STG.E.64 desc[UR6][R2.64], R4 ;
  /*0090*/  EXIT ;
  /*00a0*/  BRA 0xa0;

So if you want real 64-bit support, have fun getting your hands on a 5070! But even on sm_120, things like 64-bit × immediate 32-bit take a UIMAD.WIDE.U32 + UIMAD + UIADD3 sequence, so the support isn't all that complete.

(I've been looking into the specifics of CUDA integer arithmetic for some time now, since I've had the mad idea of doing 'horizontal' 448-bit integer arithmetic by storing one word in each thread and using the warp-shuffle instructions to send carries up and down. Given that the underlying arithmetic is all 32-bit, it doesn't make any sense to store more than 31 bits per thread. Then again, I don't know whether this mad idea makes any sense in the first place, until I implement and profile it.)

It's always been less of a big deal than it used to be -- at least on "big" GPUs -- but the article isn't really about avoiding branches. The code presented is already branchless. The people giving out the advice seem to think they are avoiding branches as optimisation but their understanding of what branching code is is apparently based on if they can see some sort of conditional construct.

This article is also relevant: https://medium.com/@jasonbooth_86226/branching-on-a-gpu-18bf...

“If you consult the internet about writing a branch of a GPU, you might think they open the gates of hell and let demons in. They will say you should avoid them at all costs, and that you can avoid them by using the ternary operator or step() and other silly math tricks. Most of this advice is outdated at best, or just plain wrong.

Let’s correct that.”

Processors change, compilers change. If you care about such details, best to ship multiple variants and pick the fastest one at runtime.

As I've mentioned here several times before, I've made code significantly faster by removing the hand-rolled assembly and replacing it with plain C or similar. While the assembly might have been faster a decade or two ago, things have changed...

Funnily enough, this is sort of what the NVIDIA drivers do: they intercept game shaders and replace them by custom ones optimized by NVIDIA. Which is why you see stuff like this in NVIDIA drivers changelog: "optimized game X, runs 40% faster"

I'd love to read more about this, what kind of changes they make and how many games they do it for. Do they ever release technical articles about it?

The other commenter makes it sound a bit more crazy than it is. "Intercept shaders" sounds super hacky, but in reality, games simply don't ship with compiled shaders. Instead they are compiled by your driver for your exact hardware. Naturally that allows the compiler to perform more or less aggressive optimisations, similar to how you might be able to optimise CPU programs by shipping C code and only compiling everything on the target machine once you know the exact feature sets.

I don't know -- if that other commenter is correct, it does sound pretty crazy.

Improving your compiler for everybody's code is one thing.

But saying, if the shader that comes in is exactly this code from this specific game, then use this specific precompiled binary, or even just apply these specific hand-tuned optimizations that aren't normally applied, that does seem pretty crazy to me.

But I don't know which it is?

Finger printing based on shaders is quite rare, really most of the time we detect things like the exe name calling us or sometime, very rarely they will give us a better name through an extension. (unreal engine does this automatically). From there all the options are simple, but full shader replacements are one. In the api I work on the shaders have a built in hash value, so that along with the game identified means we know exsactly what shader it is. Most of the replacements aren't complicated though, it's just replacing slow things with faster things for our specific hardware. In the end we are the final compiler so us tweaking things to work better should be expected to a degree.

> most of the time we detect things like the exe name calling us

What could possibly go wrong? :)

quack.exe performing noticeably worse in benchmarks on certain cards than quake.exe in the late 90's/early 2000's?

Yup. Looking at you, ATI.

For those who don't remember:

https://web.archive.org/web/20230819072628/https://techrepor...

In the case of the Minecraft mod Sodium, which replaces much of Minecraft's rendering internals, Nvidia optimisations caused the game to crash. So the mod devs had to implement workarounds to stop the driver from detecting that Minecraft is running... (changing the window title among other things)

https://github.com/CaffeineMC/sodium/issues/1486

https://github.com/CaffeineMC/sodium/issues/1816

And plain Minecraft did the opposite by adding -XX:JavaHeapDumpPath=MojangTricksIntelDriversForPerformance_javaw.exe_minecraft.exe.heapdump to the command line, when they changed the way the game started so that it wasn't detected as minecraft.exe any more. (Side effect: if you trigger a heap dump, it gets that name)

Yuuuk.

I mean this already happened [1]. But it's either that or games running like crap because they're not properly tested on our cards.

[1] https://www.neowin.net/news/yandex-alleges-amds-windows-driv...

> or games running like crap because they're not properly tested on our cards.

The decision on what game testing is proper surely lies with the game dev, not the card dev.

Plus game devs generally know that non-testing cannot cause a game to run like crap. Card devs ought to know that too.

Card devs also ought to respect that a game dev may have a good business reason for leaving his game running like crap on certain cards.

I would gauge properly testing on our card to be running the game occasionally on our cards throughout the whole development process. Currently it's been a lot of games 2 weeks out from release dropping a ton of issues on us to try and figure out. I don't even care if they don't do any optimization for our cards as long as they give us a fighting chance to figure out the issues on our side.

> Card devs also ought to respect that a game dev may have a good business reason for leaving his game running like crap on certain cards.

I can't agree with this though, business decisions getting in the way of gamers enjoying games should never be something that is settled for. If the hardware is literally to old and you just can't get it to work fine, but when a top of the line card is running below 60 fps that's just negligence.

> business decisions getting in the way of gamers enjoying games should never be something that is settled for.

That's between the game dev and his customer. None of a card dev's business.

Graphics drivers on Windows definitely do plenty of 'optimizations' for specific game executables, from replacing entire shaders to 'massaging/fixing' 3D-API calls.

And AFAIK Proton does things like this too, but for different reasons (fixing games that don't adhere to the D3D API documentation and/or obviously ignored D3D validation layer messages).

Imagine being the dev competing game Y and seeing the changelog.

It wouldn't be surprising to find out Nvidia talks directly with game developers to give them hints as to how to optimize their games

It wouldn't be surprising at all seeing as they are even in the credits for most games.

Wow, how did they pick which games to optimize?

Does the studio pay them to do it? Because Nvidia wouldn't care otherwise?

Does Nvidia do it unasked, for competitive reasons? To maximize how much faster their GPU's perform than competitors' on the same games? And therefore decide purely by game popularity?

Or is it some kinda of alliance thing between Nvidia and studios, in exchange for something like the studios optimizing for Nvidia in the first place, to further benefit Nvidia's competitive lead?

AFAIK NVIDIA and AMD do this unasked for popular game releases because it gives them a competitive advantage if 'popular game X' runs better on NVIDIA than AMD (and vice versa). If you're an AAA studio you typically also have a 'technical liason' at the GPU vendors though.

It's basically an arms race. This is also the reason why graphics drivers for Windows are so frigging big (also AFAIK).

Double "AFAIK" makes me trust you more :)

I think this is very accurate. The exception is probably those block buster games. Those probably get direct consultancy from NVIDIA during the development to make them NVIDIA-ready from day 1.

Even smaller games often get some amount of consultation from Nvidia/AMD - they are typically in the credits in some form.

I can't give details on how we do our selections (not nvidia but another gpu manufacturer). But we do have direct contacts into a lot of studios and we do try and help them fix their game if possible before ever putting something in the driver to fix it. Studios don't pay us, it's mutually benefital for us to improve the performance of the games. It also help the game run better on our cards by avoiding some of the really slow stuff.

In general if our logo is in the game, we helped them by actually writing code for them, if it's not then we might have only given them directions on how to fix issues in their game or put something in the driver to tweak how things execute. From an outside perspective (but still inside on the gpu space) nvidia does give advice to keep their competitive advantage. In my experience so far ignoring barriers that are needed as per the spec, defaulting to massive numbers when the gpu isn't known ("batman and tessellation" should be enough to find that), and then doing out right weird stuff that doesn't look like something any sane person would do in writing shaders (I have a thought in my head for that one, but it's not considered public knowledge. )

I don't work on the nvidia side of things but it's likely to be the same. Shader replacement is only one of a whole host of things we can do to make games run faster. It's actually kind of rare for use to do them since it boats the size of the driver so much. A lot of our options do change how shaders work though, like forcing a shader to use double precision floats instead of the single it was compiled with.

> A lot of our options do change how shaders work though, like forcing a shader to use double precision floats instead of the single it was compiled with.

What benefit would that give? Is double precision faster than single on modern hardware?

I would guess it's done for compatibility or quality reasons

Only more precision. But no, doubles are not faster. At best they’re the same instruction latency & throughput as singles, and that’s only on a few expensive pro/datacenter GPUs. Even if they are technically the same instruction speed, they’re still 2x the memory & register usage, which can compromise perf in other ways. Doubles on consumer GPUs are typically anywhere from 16 to 64 times slower than singles.

FWIW, I’ve never heard of shader replacement to force doubles. It’d be interesting to hear when that’s been used and why, and surprising to me if it was ever done for a popular game.

That's specifically because gpus aren't IEEE compliant, and calculations will drift differently on different gpus. Double precision can help avoid divide by zero errors in some shaders because most don't guard against that and NANs propagate easily and show up as visual corruption.

What GPUs aren’t IEEE compliant? I don’t think that’s been true for more than a decade.

After a bunch of testing and looking around I think I should actually change my statement. GPUs do offer IEEE floating point compliance by default, but don't strictly adhere to it. Multiple optimizations that can be applied by the driver developers can massively effect the floating point accuracy.

This is all kind of on the assumption that the accuracy of floating point multiplication and division is in the IEEE spec, I was told before that it was but searching now I can't seem to find it one way or the other.

I believe one of the optimizations done by nvidia is to drop f32 variables down to f16 in a shader. Which would technically break the accuracy requirement (as before if it exists). I don't have anything I can offer as proof of that due to NDA sadly though. I will note that most of my testing and work is done in PIX for Windows, and most don't have anti-cheat so they're easy to capture.

What shaders (presumably GLSL & HLSL) do precision wise isn’t an IEEE compliance issue, it’s either a DX/Vulkan spec issue, OR a user compiler settings issue. Dropping compliance is and should be allowed when the code asks for it. This is why GLSL has lowp, mediump, and highp settings. I think all GPUs are IEEE compliant and have been for a long time.

I agree on the dropping compliance when asked for aspect, the problem I'm referring to more is the driver dropping compliance without the game asking for it. If the underlying system can randomly drop compliance when ever it thinks it's fine without telling the user and without the user asking, I would not consider that compliant.

That is fair, if true. But I’m very skeptical any modern drivers are dropping compliance without being asked. The one possibility I could buy is that you ran into a case of someone having dropped precision intentionally in a specific game in order to “optimize” a shader. Otherwise, precision and IEEE compliance is the prerogative of the compiler. The compiler is sometimes in the driver, but it never decides on it’s own what precision to use, it uses either default or explicit precision settings. The only reason it would not produce IEEE compliant code is if it was being asked to.

> > A lot of our options do change how shaders work though, like forcing a shader to use double precision floats instead of the single it was compiled with.

That will break code sufficienly reliant on the behaviour of sungle precision, though.

Obviously, which is the reason you don't do something like that without appropriate amounts of testing.

Testing on platforms that perverts the precision is outside appropriate, I would say.

Code depending on specific low precision is very likely to be so fragile it won't make it anywhere near release.

Consider an RNG.

I consider a floating point RNG in a shader that is seeded for reproducibility to be a bad idea.

Why, pray?

GPUs often have weird sizes for optimal behavior and you have to depend on math-based optimizations never being applied to your code.

In my opinion, floating point shaders should be treated as a land of approximations.

You asked in another comment why /width*width isn't optimized out by the compiler. But it's changes just like that that will break an RNG!

> In my opinion, floating point shaders should be treated as a land of approximations.

Fine, but that leaves you responsible for the breakage the shader of an author that holds the opposite opinion, as he is entitled to do. Precision =/= accuracy.

But do you want /width*width to be optimized? Or associative and commutative operations in general? Then you have to reject the opposite opinion.

> But do you want /width*width to be optimized?

I do. But I don't see such optimisation as anything to with your changing float type.

Changes like that will break just as much code as adding extra precision will. Because it will change how things round, and not much else, just like adding extra precision. They're both slightly disruptive, and they tend to disrupt the same kind of thing, unlike removing precision which is very disruptive all over.

"Fragile shader code won't make it anywhwere near release" is a pretty bold claim.

If it happens to work with the Nvidia driver it's getting shipped.

But since the premise is that any difference breaks the code, will it work the exact same way across different devices? That's what I'm skeptical of.

Unless every dev and tester that touches this shader is using the same hardware, which seems like an obvious mistake to avoid...

I will note, half of the customer facing bugs I get are "works on nvidia." Only to find out that it is a problem with the game and not the driver. Nvidia allows you to ignore a lot of the spec and it causes game devs to miss a lot of obvious bugs. A few examples:

1) Nvidia allows you to write to read only textures, game devs will forget to transition them to writable and will appear as corruption on other cards.

2) Nvidia automatically work with diverging texture reads, so devs will forget to mark them as a nonuniform resource index, which shows up as corruption on other cards.

3) Floating point calculations aren't IEEE compliant, one bug I fixed was x/width*width != x, On Nvidia this ends up a little higher and on our cards a little lower. The game this happened on ended up flooring that value and doing a texture read, which as you can guess, showed up as corruption on our cards.

1 and 2 are specifically required by the microsoft directx 12 spec, but most game devs aren't reading that and bugs creep in. 3 is a difference in how the ALU is designed, our cards being a little closer to IEEE compliant. A lot of these issue are related to how the hardware works, so stays pretty consistent between the different gpus of a manufacturer.

Side note: I don't blame the devs for #3, the corruption was super minor and the full calculation was spread across multiple functions (assumed by reading the dxil). The only reason it sticks out in my brain though is because the game devs were legally unable to ever update the game again, so I had to fix it driver side. That game was also Nvidia sponsored, so it's likely our cards weren't tested till very late into the development. (I got the ticket a week before the game was to release.) That is all I'm willing to say on that, I don't want to get myself in trouble.

> Floating point calculations aren't IEEE compliant

To late to edit, but I want to half retract this statement, they are IEEE compliant, but due to optimizations that can be applied by the driver developers they aren't guaranteed to be. This is assuming that the accuracy of a multiply and divide are specified in the IEEE floating point spec, I'm seeing hints that it is, but I can't find anything concrete.

> the game devs were legally unable to ever update the game again, so I had to fix it driver side.

Could they not intercept the calls to inject a fix?

I'm just going off what I was told there, I was forced to make the fix since the game developers no longer were partnered to the company that owned the license to the content.

PS Why is /width*width not eliminated on compilation?

Good question, I'm assuming it's due to the calculation happen across a memory barrier of some kind or due to all the branches in between so llvm is probably avoiding the optimization. It was quite a while ago so it is something I could re-investigate and actually try and fix. I would have to wait for downtime with all the other tickets I'm getting though. It's also just something that dxc itself should be doing, but I have no control over that.

> But since the premise is that any difference breaks the code,

Not any.

> will it work the exact same way across different devices

Yes, where they run IEEE FP format.

> Not any.

Just about any. It's pretty difficult to write code where changing the rounding of the last couple bits breaks it (as happens if you use wider types during the calculation), but other changes don't break it.

What real code have you seen with that behavior?

> but other changes don't break it.

Why should that be a requirement? Obviously the driver can make other changes that can break any code.

> Why should that be a requirement?

Originally I said "the premise is that any difference breaks the code".

You replied with "Not any."

That is where the requirement comes from, your own words. This is your scenario, and you said not all differences would break the hypothetical code.

This is your choice. Are we talking about code where any change breaks it (like a seeded/reproducible RNG), or are we talking about code where there are minor changes that don't break it but using extra precision breaks it? (I expect this category to be super duper rare)

> Originally I said "the premise is that any difference breaks the code".

> You replied with "Not any."

> That is where the requirement comes from, your own words.

I wasn't stating a requirement. I was disputing your report of the premise.

> or are we talking about code where there are minor changes that don't break it but using extra precision breaks it?

Yup.

> > or are we talking about code where there are minor changes that don't break it but using extra precision breaks it?

> Yup

Then I will ask again, have you ever seen code that falls into this category? I haven't. And an RNG would not fall into this category.

In the case that does happen, then we don't apply that setting. Most of the changes applied are extensively tested and toggles like that are more often used for already broken shaders.

> In the case that does happen

You meant "in the case we find where that does happen"?

Little consolation for the cases you don't find because e.g. you cannot afford to play a game to the end.

Fair enough, I can't say anything I've done has ever caused an issue like that (A new ticket would have been made and sent to me.) But I also can't say that it has never happened, so I'm not really in a position to disagree. We do have a good QA team though and we have an "open" beta program that also catches a lot of issues before they become more widely public.

I for one would be most unhappy to find my game failing on card X due to drivers perverting my float type.

> "optimized game X, runs 40% faster"

... and looks 4O% crappier? E.g. stuttery, because the driver does not get to see the code ahead of time.

I think figuring out the fastest version of a shader at runtime is very non-trivial, I'm not aware of any game or engine that can do this.

I think it'd be possible in principle, because most APIs (D3D, GL, Vulkan etc) expose performance counters (which may or may not be reliable depending on the vendor), and you could in principle construct a representative test scene that you replay a couple times to measure different optimizations. But a lot of games are quite dynamic, having dynamically generated scenes and also dynamically generated shaders, so the number of combinations you might have to test seems like an obstacle. Also you might have to ask the user to spend time waiting on the benchmark to finish.

You could probably just do this ahead of time with a bunch of different GPU generations from each vendor if you have the hardware, and then hard-code the most important decision. So not saying it'd be impossible, but yeah I'm not aware of any existing infrastructure for this.

The last time I did anything like this (it was for CPU linear algebra code designed to run in very heterogeneous clusters), I first came up with a parameterization that approximated how I'd expect an algorithm to perform. Then, once for each hardware combination, you sweep through the possible parameterization space. I used log-scaled quantization to make it cheap to index into an array of function pointers based on input specifics.

The important thing to note is that you can do that computation just once, like when you install the game, and it isn't that slow. Your parameterization won't be perfect, but it's not bad to create routines that are much faster than any one implementation on nearly every architecture.

you'd only have to test worst/median case scenes, which you could find with a bit of profiling!

This would be acceptable if it meant adding one more shader, but with "modern" graphics APIs forcing us to sometimes have thousands of permutations for the same shader, every variant you add multiplies that count by 2x.

We also don't have an infinite amount of time to work on each shader. You profile on the hardware you care about, and if the choice you've made is slower on some imaginary future processor, so be it - hopefully that processor is faster enough that this doesn't matter.

Graphics APIs don't force your to have thousands of shaders. The abstraction in your engine might.

Some of the mistakes/confusion being pointed out in the article is being replicated here, it seems.

The article is not claiming that conditional branches are free. In fact, the article is not making any point about the performance cost of branching code, as far as I can tell.

The article is pointing out that conditional logic in the form presented does not get compiled into conditionally branching code. And that people should not continue to propagate the harmful advice to cover up every conditional thing in sight[0].

Finally, on actually branching code: that branching code is more complicated to execute is self-evident. There are no free branches. Avoiding branches is likely (within reason) to make any code run faster. Luckily[1], the original code was already branchless. As always, there is no universal metric to say whether optimisation is worthwhile.

[0] the "in sight" is important -- there's no interest in the generated code, just in the source code not appearing to include conditional anythings.

[1] No luck involved, of course ... (I assume people wrote to IQ to suggest apparently glaringly obvious (and wrong) improvements to their shader code, lol)

So why isn't the compiler smart enough to see that the 'optimised' version is the same?

Surely it understands "step()" and can optimize the "step()=0.0" and "step()==1.0" cases separately?

This is presumably always worth it, because you would at least remove one multiplication (usually turning it into a conditional load/store/something else)

It may very well be, it is the type of optimisation where it is quite possible that some compilers may do it some of the time, but it is definitely also possible to write a version that the compiler can't grok.

Yup, they most likely do. After all everything is LLVM based nowadays.

That's not true for shader compilers included in drivers - some use LLVM but definitely not all of them.

The other part of the optimization issue is that you can't take to long to try anything and everything. Most of the optimizations happen on the driver side, and anything that takes to long will show up as shader compilation stutter. I can't say currently if this is or isn't done, it's just always something you have to think about.

Thanks Inigo !

  The second wrong thing with the supposedly optimizer version is that it actually runs much slower than the original version. The reason is that the step() function is actually implemented like this:

  float step( float x, float y )
  {
      return x < y ? 1.0 : 0.0;
  }

How are we supposed to know what OpenGL functions are emulated rather than calling GPU primitives ?

The only way is do what OP did – compile your shader, disassemble, and read the assembly.

I do that quite often with my HLSL shaders, learned a lot about that virtual instruction set. For example, it’s interesting GPUs have instruction sincos, but inverse trigonometry is emulated while compiling.

This is a great question that I see everywhere in programming and I think it is core to why you measure first when optimizing.

You generally shouldn’t know or care how a built in is implemented. If do care, you’re probably thinking about optimization. At that point the answer is “measure and find out what works better.”

EDIT: I see my source of confusion must be that "branch" must have a well-understood hardware-specific meaning that goes beyond the meaning I grew up with, which is that a conditional is a branch, because the path control takes (at the machine code level) is chosen at runtime. This makes a conditional jump a branch by definition.

> How are we supposed to know what OpenGL functions are emulated rather than calling GPU primitives?

To me the problem was obvious, but then again I'm having trouble with both your and author's statements about it.

The problem I saw was, obviously by going for a step() function, people aren't turning logic into arithmetic, they're just hiding logic in a library function call. Just because step() is a built-in or something you'd find used in mathematical paper doesn't mean anything; the definition of step() in mathematics is literally a conditional too.

Now, the way to optimize it properly to have no conditionals, is you have to take a continuous function that resembles your desired outcome (which in the problem in question isn't step() but the thing it was used for!), and tune its parameters to get as close as it can to your target. I.e. typically you'd pick some polynomial and run the standard iterative approximation on it. Then you'd just have an f(x) that has no branching, just a bunch of extra additions and multiplications and some "weirdly specific" constants.

Where I don't get the author is in insisting that conditional move isn't "branching". I don't see how that would be except in some special cases, where lack of branching is well-known but very special implementation detail - like where the author says:

> also note that the abs() call does not become a GPU instruction and instead becomes an instruction modifier, which is free.

That's because we standardized on two's complement representation for ints, which has the convenient quality of isolating sign as the most significant bit, and for floats the representation (IEEE-754) was just straight up designed to achieve the same. So in both cases, abs() boils down to unconditionally setting the most significant bit to 0 - or, equivalently, masking it off for the instruction that's reading it.

step() isn't like that, nor any other arbitrary ternary operation construct, and nor is - as far as I know - a conditional move instruction.

As for where I don't get 'ttoinou:

> How are we supposed to know what OpenGL functions are emulated rather than calling GPU primitives

The basics like abs() and sqrt() and basic trigonometry are standard knowledge, the rest... does it even matter? step() obviously has to branch somewhere; whether you do it yourself, let a library do it, or let the hardware do it, shouldn't change the fundamental nature.

It kinda does when you’re wondering what’s going on in backstage and working with shaders on multiple OS, drivers and hardware.

   Now, the way to optimize it properly to have no conditionals, is you have to take a continuous

I suspect that we shaders authors really like Clean Math and that’s also why we like to think such “optimizations” with the step function is a nice modification :-)

A "branch" here is a conditional jump. This has the issues the article mentions, which branchless programming avoids:

    Again, there is no branching - the instruction pointer isn't manipulated, there's no branch prediction involved, no instruction cache to invalidate, no nothing.

This has nothing to do with whether the behaviour of some instruction depends on its arguments. Looking at the Microsoft compiler output from the article, the iadd (signed add) instruction will get different results depending on its arguments, and the movc (conditional move) will store different values depending on its arguments, but after each the instruction pointer will just move onto the next instruction, so there are no branches.

Branching is different instruction paths, so it requires reading the instructions from different memory that causes a delay jumping to those new instructions rather than plowing ahead on the current stream of instructions. So a conditional jump is a branch but a conditional move is just an instruction that moves one of two values into a register but doesn’t affect what code is executed next.

> the meaning I grew up with, which is that a conditional is a branch

A conditional jump is a branch. But a branch has always had a different meaning than a generic “conditional”. There are conditional instructions that don’t jump, e.g. CMP, and the distinction is very important. Branch or conditional jump means the PC can be set to something other than ‘next instruction’. A conditional, such a conditional select or conditional move, one that doesn’t change the PC, is not a branch.

> take a continuous function […] Then you’d just have an f(x) that has no branching

One can easily implement conditional functions without branching. You can use a compare instruction followed by a Heaviside function on the result, evaluate both sides of the result, and sum it up with a 2D dot product (against the compare result and its negation). That is occasionally (but certainly not always) faster on a GPU than using if/else, but only if the compiler is otherwise going to produce real branch instructions.

Maybe I’m misunderstanding why branching is slow on a GPU. My understanding was that it’s because both sides of the branch are always executed, just one is masked out (I know the exact mechanics of this have changed), so that the different cores in the group can use the same program counter. Something to that effect, at least.

But in this case, would calculating both sides and then using a way to conditionally set the result not perform the same amount of work? Whether you’re calculating the result or the core masks the instructions out, it’s executing instructions for both sides of the branch in both cases, right?

On a CPU, the performance killer is often branch prediction and caches, but on the GPU itself executing a mostly linear set of instructions, or is my understanding completely off? I guess I don’t really understand what it’s doing, especially for loops.

The primary concern is usually over the masking you’re talking about, the issue being simply that you’re proportionally cutting down the number of threads doing useful work. Using Nvidia terminology, if only one thread in a warp is active during a branch, the GPU throughput is 32x slower than it could be with a full warp.

Not all GPU branches are compiled in a straight line without jumps, so branching on a GPU does sometimes share the same instruction cache churn that the CPU has. That might be less of a big deal than thread masking, but GPU stalls still take dozens of cycles. And GPUs waiting on memory loads, whether it’s to fill the icache or anything else, are up to 32x more costly than CPU stalls, since all threads in the warp stall.

Loops are just normal branches, if they’re not unrolled. The biggest question with a loop is will all threads repeat the same number of times, because if not, the threads that exit the loop have to wait until the last thread is done. You can imagine what that might do for perf if there’s a small number of long-tail threads.

> a conditional is a branch, because the path control takes (at the machine code level) is chosen at runtime. This makes a conditional jump a branch by definition.

s/conditional is a/conditional jump is a/

Problem solved.

Non-jump conditionals have been a thing for decades.

I’ve never seen a GPU with special primitives for any functions than you’d see in pc style assembly. Every time I’ve looked at a decompiled shader, it’s always been pretty much what you think of in C.

Aldo specs like OpenGL specify many intrinsic behavior, which is then implemented as the spec, using standard assembly instructions.

Find an online site that decompiles to various architectures.

Why are you supposed to know?

Because you care about performance? step being implemented as a libray function on top of a conditional doesn't really say anything about its performance vs being a dedicated instruction. Don't worry about the implementation.

Because you are curious about GPU architectures? Look at disassembly, (open source) driver code (including LLVM) and/or ISA documentation.

I've been caught by this. Even Claude/ChatGPT will suggest it as an optimisation. Every time I've measured a performance drop doing this. Sometimes significant.

Is that weird? LLMs will just repeat what is in their training corpus. If most of the internet is recommending something wrong (like this conditional move "optimization") then that is what they will recommend too.

Not weird but important to note.

> Even Claude/ChatGPT will suggest it as an optimisation.

LLMs just repeat what people on the internet say, and people are often wrong.

Hmm, godbolt is showing branches in the vulkan output:

        return x>0.923880?vec2(s.x,0.0):
               x>0.382683?s*sqrt(0.5):
                          vec2(0.0,s.y);

turns into

         %24 = OpLoad %float %x
         %27 = OpFOrdGreaterThan %bool %24 %float_0_923879981
               OpSelectionMerge %30 None
               OpBranchConditional %27 %29 %35
         %29 = OpLabel
         %31 = OpAccessChain %_ptr_Function_float %s %uint_0
         %32 = OpLoad %float %31
         %34 = OpCompositeConstruct %v2float %32 %float_0
               OpStore %28 %34
               OpBranch %30
         %35 = OpLabel
         %36 = OpLoad %float %x
         %38 = OpFOrdGreaterThan %bool %36 %float_0_382683009
               OpSelectionMerge %41 None
               OpBranchConditional %38 %40 %45
         %40 = OpLabel
         %42 = OpLoad %v2float %s
         %44 = OpVectorTimesScalar %v2float %42 %float_0_707106769
               OpStore %39 %44
               OpBranch %41
         %45 = OpLabel
         %47 = OpAccessChain %_ptr_Function_float %s %uint_1
         %48 = OpLoad %float %47
         %49 = OpCompositeConstruct %v2float %float_0 %48
               OpStore %39 %49
               OpBranch %41
         %41 = OpLabel
         %50 = OpLoad %v2float %39
               OpStore %28 %50
               OpBranch %30
         %30 = OpLabel
         %51 = OpLoad %v2float %28
               OpReturnValue %51

https://godbolt.org/z/aqob7YfWq

Vulcan opcode shader lang is not executed. It’s a platform neutral intermediate language, so won’t have the special purpose optional instructions most GPUs do since GPUs aren’t required to.

It likely compiles down on the relevant platforms as the original article did.

Unrelated, but somehow similar: I really hate it that it's not possible to force gcc to transform things like this into a conditional move:

x > c ? y : 0.;

It annoyed me many times and it still does.

What do you mean? Do you want to annotate the condition as unpredictable, so that the compiler always assumes that a conditional move is beneficial?

(Compilers obviously do this transformation, including GCC, but it is not always beneficial, especially on x86-64.)

And it's not always possible! E.g. most RISC-V CPUs don't support it yet.

Eh, it takes ~3-4 instrs to do a branchless "x ? y : z" on baseline rv64i (depending on the format you have the condition in) via "y^((y^z)&x)", and with Zicond that only goes down to 3 instrs (they really don't want to standardize GPR instrs with 3 operands so what Zicond adds is "x ? y : 0" and "x ? 0 : y" ¯\_(ツ)_/¯; might bring the latency down by an instr or two though).

It's more about removing branches than instruction counts or latency.

The "y^((y^z)&x)" method is already branchless, and close in performance to the Zicond variant, is my point; i.e. Zicond doesn't actually add much.

Are you sure? As soon as you add actual computations in you're heading through the whole execution pipeline & forwarding network, tying up ALUs, etc. Zicond can probably be handled without all that.

Also that isn't actually equivalent since `x` needs to be all 1s or all 0s surely? Neither GCC nor Clang use that method, but they do use Zicond.

Zicond's czero.eqz & czero.nez (& the `or` to merge those together for the 3-instr impl of the general `x?y:z`) still have to go through the execution pipeline, forwarding network, an ALU, etc just as much as an xor or and need to. It's just that there's a shorter dependency chain and maybe one less instr.

Indeed you may need to negate `x` if you have only the LSB set in it; hence "3-4 instrs ... depending on the format you have the condition in" in my original message.

I assume gcc & clang just haven't bothered considering the branchless baseline impl, rather than it being particularly bad.

Note that there's another way some RISC-V hardware supports doing branchless conditional stores - a jump over a move instr (or in some cases, even some arithmetic instructions), which they internally convert to a branchless update.

Yes, that would be great. It's not always benefical, but in some (rare, but for me important) cases it's better. Currently, the only way to ensure a conditional move is used, is to use inline assembly. This is not portable and also less maintainable than a "proper" solution.

clang has the __builtin_unpredictable() intrinsic[0] for this purpose.

[0] https://clang.llvm.org/docs/LanguageExtensions.html#builtin-...

Seems to work just fine on gcc and clang?

https://www.godbolt.org/z/ffEvvjhz8

PS: and it also doesn't matter whether a ternary is used or a traditional if (as one would expect):

https://www.godbolt.org/z/zjb4KdqvK

(the float version also appears to not use branches: https://www.godbolt.org/z/98bdheKK4)

For such simple expression I would expect the compiler to pick the right output pattern based on the target CPU though...

Not always - https://www.godbolt.org/z/zYxeahf3T. And for any modern (as in, made in the last two decades) x86 processor the branchless version will be hilariously better if the condition is unpredictable (which is a thing the compiler can't know by itself, hence wanting to have an explicit way to request a conditional move instr) and the per-branch code takes less than like multiple dozens of cycles.

Worse, doing one of the idioms for a conditional move ends up getting gcc to actually produce a conditional move, but clang doesn't, even with its __builtin_unpredictable: https://www.godbolt.org/z/bq9axzvjG

You want to pass -mllvm -x86-cmov-converter=false. I assume LLVM has a pass to undo conditional moves on x86 whenever a heuristic determines skipping a calculation by branching is cheaper than doing the calculation and using a conditional move.

Unfortunately, the heuristic that calculates the expense often gets things wrong. That is why OpenZFS passes -mllvm -x86-cmov-converter=false to Clang for certain files where the LLVM heuristic was found to do the wrong thing:

https://github.com/openzfs/zfs/commit/677c6f8457943fe5b56d7a...

There is an open LLVM issue regarding this:

https://github.com/llvm/llvm-project/issues/62790

The issue explains why __builtin_unpredictable() does not address the problem. In short, the metadata is dropped when an intermediate representation is generated inside clang since the IR does not have a way to preserve the information.

Yeah, the need of the "-mllvm -x86-cmov-converter=false" hack is stupid; forgot to check with it. In my mind I guess I equivocated that being fixed with https://reviews.llvm.org/D118118 (and indeed __builtin_unpredictable does work with ?:), but, no, that flag still improves things today for the rest of the cases.

It just occurred to me that Clang and GCC did not necessarily fail to use conditional moves in your examples. While they failed to use explicit cmov instructions, cmp/jmp 1 instruction/mov is actually an idiom for an implicit cmov. Some CPU instruction decoders are able to turn it into a cmov without an explicit cmov instruction. In the case of RISC-V, the designers are philosophically opposed to explicit cmov instructions and expect compilers to generate this idiom for CPUs that support cmov. I asked them to implement cmov virtual instructions to be nice to people reading RISC-V assembly, but I am not sure if anything will come of it:

https://github.com/riscv/riscv-bitmanip/issues/185

I do not know if any x86 CPUs recognize the implicit cmov idiom offhand, but if any do, then while an extra instruction was used, the conditional move would still be done on those that recognize the idiom.

By the way, I just noticed a case where you really don’t want the compiler to generate a cmov, explicit or otherwise since it would risk division by zero:

https://github.com/openzfs/zfs/commit/f47f6a055d0c282593fe70...

Here is a godbolt link showing some output:

https://www.godbolt.org/z/4daKTKqfr

Interestingly, Clang correctly does not generate a cmov (implicit or explicit) for the outer ternary operation, while it does generate an explicit cmov for the inner ternary operator in MIN() without -mllvm -x86-cmov-converter=false. Passing -mllvm -x86-cmov-converter=false to Clang does not change the output, which makes Clang’s behavior correct.

GCC will not generate cmov for either ternary operator, which while also technically correct, is slow. This could still have been an implicit conditional move had GCC not avoided the implicit cmov idiom.

Using GCC’s __builtin_expect_with_probability() in MIN() does not cause GCC to change its output. If we remove the outer ternary, GCC will happily generate a cmov instruction. Given that GCC generally assumes that undefined behavior is not invoked to make code faster and will happily generate the cmov when there is a division by 0 bug, it is odd that upon seeing a check that verifies the assumption GCC made is true, GCC decides to stop generating a cmov. I am sure the way GCC does things is much more complicated than my interpretation of the output, but the behavior is odd enough to merit a comment.

I haven't heard of anything outside RISC-V having jump-over-mov as an idiom (though I've heard of potentially some CPUs having the ability to unwind only necessary parts on mispredictions over small bits of code or something; still some misprediction penalty though I believe; and even with -march=haswell behavior doesn't change).

I find the RISC-V solution (which fwiw I mentioned in a sibling thread[0]) rather sad; there's no way to check whether it's implemented, and even where it is I could imagine it being problematic (i.e. if the instructions cross a fetch block or cacheline or something and it gets ran as a branch, or some instrs around it break the fusion pattern checking), and where it's unsupported or otherwise doesn't work properly it'll "work" but be horrifically slow.

fwiw I haven't ever seen __builtin_expect_with_probability actually do anything for unpredictable branches; I just included it in my compiler explorer link for completeness.

Using a version of MIN that caches the X/Y computations gets gcc to produce a cmov, but makes clang's output longer: https://www.godbolt.org/z/6h8obxKG8

[0]: https://news.ycombinator.com/item?id=42992533

You might want to give feedback to the risc-v developers (although it might be too late at this point). What is the way to check if implicit cmov instructions are implemented in the CPU instruction decoder?

If AMD did not implement this in Zen 5, maybe we could ask them to add it in Zen 7 or 8. I assume it would be too late to ask them to add this in Zen 6.

Thanks for the caching tip.

There's of course no "way" to check, as it's a microarchitectural property. Your best bet is comparing performance of the same code on predictable vs unpredictable branches.

I don't think there's any need for x86 cores to try to handle this; it's just a waste of silicon for something doable in one instruction anyway (I'd imagine that additionally instruction fusion is a pretty hot path, especially with jumps involved; and you'll get into situations of conflicting fusions as currently cmp+jcc is fused, so there's the question of whether cmp+jcc+mov becomes (cmp+jcc)+mov or cmp+(jcc+mov), or if you have a massive three-instruction four-input(?) fusion).

Oh, another thing I don't like about fusing condjump+mv - it makes it stupidly more non-trivial to intentionally use branches on known-predictable conditions for avoiding the dependency on both branches.

> There's of course no "way" to check, as it's a microarchitectural property. Your best bet is comparing performance of the same code on predictable vs unpredictable branches.

I was afraid the answer to my question would be that, but since my read of your previous comment “there's way to check whether it's implemented” seemed to suggest you knew a way I did not, I had my fingers crossed. At least, it had been either that you knew a trick I did not, or that a typo had deleted the word “no”.

> I don't think there's any need for x86 cores to try to handle this; it's just a waste of silicon for something doable in one instruction anyway (I'd imagine that additionally instruction fusion is a pretty hot path, especially with jumps involved; and you'll get into situations of conflicting fusions as currently cmp+jcc is fused, so there's the question of whether cmp+jcc+mov becomes (cmp+jcc)+mov or cmp+(jcc+mov), or if you have a massive three-instruction four-input(?) fusion).

Interestingly, the RISC-V guys seem to think that adding an explicit instruction is a waste of silicon while adding logic to detect the idiom to the instruction decoder is the way to go. x86 cores spend enormous amounts of silicon on situational tricks to make code run faster. I doubt spending silicon on one more trick would be terrible, especially since the a number of other tricks to extract more performance from things likely apply to even more obscure situations. As for what happens in the x86 core, the instruction decoder would presumably emit what it emits for the explicit version when it sees the implicit version. I have no idea what that is inside a x86 core. I suspect that there are some corner cases involving the mov instruction causing a fault to handle (as you would want the cpu to report that the mov triggered the fault, not the jmp), but it seems doable given that they already had to handle instruction faults in other cases of fusion.

Also, if either of us were sufficiently motivated, we might be able to get GCC to generate better code through a plugin that will detect the implicit cmov idiom and replace it with an explicit cmov:

https://gcc.gnu.org/onlinedocs/gccint/Plugins.html

A similar plugin likely could be written for LLVM:

https://llvm.org/docs/WritingAnLLVMNewPMPass.html#registerin...

Note that I have not confirmed whether their plugins are able to hook the compiler backend where they would need to hook to do this.

Of course, such plugins won’t do anything for all of the existing binaries that have the implicit idiom or any new binaries built without the plugins, but they could at least raise awareness of the issue. It is not a full solution since compilers don’t emit the implicit cmov idiom in all cases where a cmov would be beneficial, but it would at least address the cases where they do.

> since my read of your previous comment seemed to suggest you knew a way I did not, I had my fingers crossed.

Whoops, typo! edited.

> Interestingly, the RISC-V guys seem to think that adding an explicit instruction is a waste of silicon while adding this to the instruction decoder is the way to go

From what I've read, the thing they're against (or at least is a major blocker) is having a standard GPR instruction that takes 3 operands, as all current GPR instrs take a max of two. I cannot imagine there being any way that fusing instructions is less silicon than a new instruction whatsoever; if anything else, it'd be not wanting to waste opcode space, or being fine with the branchy version (which I'm not).

Zen 4, at least as per Agner's microarchitecture optimization guide, only fuses nops and cmp/test/basic_arith+jcc; not that many tricks, only quite necessary ones (nops being present in code alignment, and branches, well, being basically mandatory every couple instructions).

No need for a plugin; it is possible to achieve branchess moves on both as-is: https://www.godbolt.org/z/eojqMseqs. A plugin wouldn't be any more stable than that mess. (also, huh, __builtin_expect_with_probability actually helped there!)

I'd imagine a major problem for the basic impls is that the compiler may early on lose the info that the load can be ran in both cases, at which point doing it unconditionally would be an incorrect transformation.

I had suggested the virtual instructions to the RISC-V developers to eliminate branchy cmov assembly, as I am not happy with it either. It is surprising to realize that x86 cores are not making more use of macro-ops fusion, contrary to my expectation, but I guess it makes sense now that I think about it. Their designers have plenty of other knobs for tuning performance and the better their branch predictor becomes, the less this actually matters outside of the cases where developers go out of their way to use cmov.

A plugin would handle cases where the implicit idiom is emitted without needing the developer to explicitly try to force this. As far as I know, most people don’t ever touch conditional moves on the CPU and the few that do (myself included), only bother with it for extremely hot code paths, which leaves some dangling fruit on the table, particularly when the compiler is emitting the implicit version by coincidence. The safety of the transformation as a last pass in the compiler backend is not an issue since the output would be no more buggy than it previously was (as both branches are already calculated). Trying to handle all cases (the non-low dangling fruit) is where you have to worry about incorrect transformations.

Ah, your gcc example does have the branchful branch that could be done branchlessly; I was thinking about my original example with a load, which can't be transformed back.

On fusion, https://dougallj.github.io/applecpu/firestorm.html mentions ones that Apple's M1 does - arith+branch, and very specialized stuff.

I think the core problem is that when writing code like this you need experience be sure that it won’t have a conditional branch. How many operations past the conditional cause a branch? Which operations can the compiler elide to bring the total below this count? I’m all for writing direct code and relying on smart compilers but it’s often hard to know if and where I’m going to get bitten. Do I always have to inspect the assembly? Do I need a performance testing suit to check for accidental regressions? I find it much easier if I can give the compiler a hint on what I expect it to do, this would be similar to a @tailcall annotation. That way I can explore the design space without worry that I’ll accidentally overstep a some hard to reason about boundary that will tank the performance.

This should be quantified and generalized for a full set of cases - that way the argument would stand far more clearly.

Something like this:

https://doliveira4.github.io/gpuconditionals/

(no warranty)

It's funny because I rarely seen this (wrong approach) done anywhere else but I pick it up by myself (like a lot did I presume) and still am the first to do it everytime I see the occasion, not so for optimizations (while I admit I thought it wouldn't hurt) but for the flow and natural look of it. It feels somehow more right to me to compose effect by signals interpolations rather than clear ternary branch instructions.

Now I'll have to change my ways in fear of being rejected socially for this newly approved bad practice.

At least in WebGPU's WGSL we have the `select` instruction that does that ternary operation hidden as a method, so there is that.

I believe the conclusion is correct in 2025, but the article in a way just perpetuates the 'misinformation', making it seem like finding if your code will compile to a dynamic branch or not is easier than it is.

The unfortunate truth with shaders is that they are compiled by the users machine at the point of use. So compiling it on just your machine isn't nearly good enough. NVIDIA pricing means large numbers of customers are running 10 year old hardware. Depending on target market you might even want the code to run on 10 year old integrated graphics.

Does 10 year old integrated graphics across the range of drivers people actually have running prefer conditional moves over more arithmetic ops.. probably, but I would want to keep both versions around and test on real user hardware if this shader was used a lot.

    So, if you ever see somebody proposing this

    float a = mix( b, c, step( y, x ) );

The author seems unaware of

    float a = mix( b, c, y > x );

which encodes the desired behavior and also works for vectors:

    The variants of mix where a is genBType select which vector each returned component comes from. For a component of a that is false, the corresponding component of x is returned. For a component of a that is true, the corresponding component of y is returned.

The author doesn't seem to say that mix should be avoided. Just that you shouldn't replace a ternary with step+mix. In your quote, you left out the 2nd half of the sentence: "as an optimization to [ternary]".

The author frames his post to be about education:

    please correct them for me. The misinformation has been around for 20 years

But his education will fail as soon as you're operating on more than scalars. It might in fact do more harm than good since it leads the uneducated to believe that mix is not the right tool to choose between two values.

If you only pass a boolean 0 or 1 for the “a” mix parameter, when is using mix better than a ternary? Can you give an example? I’m not sure mix is ever the right tool to choose between two values. It’s a great tool for blending two values, for linear interpolation when “a” is between 0 and 1. But if “a” is only 0 or 1, I don’t think mix will help you, and it could potentially hurt if the two values you mix are expensive function calls.

a can be a vector of booleans.

Does ternary work for that? Inigo is saying "if ternary is working, leave it as ternary". I think you're talking about a situation where ternary wouldn't work, so Inigo isn't saying anything about that situation.

So what? That’s a language convenience, not a performance feature. Calling mix with a vec4 is no faster than using 4 scalar mix calls, and possibly much slower than using 4 ternaries. If your x & y params are expensive functions, and you use a vector of booleans, then you might be really shooting yourself in the foot.

IQ is most definitely aware of that. Are you unaware of who IQ is?

I'm going to go on a limb and claim the author is aware:

https://www.shadertoy.com/user/iq

He also is one of the main authors of the whole site... legend.

Hah, yeah, this is a brilliant "do you even know who you're talking to?" moment.

WGSL has select https://www.w3.org/TR/WGSL/#select-builtin

I don’t know enough about these implementations to know if this can be interpreted as a blanket ‘conditionals are fine’ or, rather, ‘ternary operations which select between two themselves non-branching expressions are fine’.

Like does this apply if one of the two branches of a conditional is computationally much more expensive? My (very shallow) understanding was that having, eg, a return statement on one branch and a bunch of work on the other would hamstring the GPU’s ability to optimize execution.

A real branch is useful if you can realistically skip a bunch of work, but this requires all the lanes to agree, on a GPU that means 32 to 64 lanes need to all agree, also for something basic like a few arithmetic ops there is no point.

A GPU/SIMT branch works by running both sides, unless all threads in the thread group (warp/wavefront) make the same branch decision. As long as both paths have at least one thread, the GPU will run both paths sequentially and simply set the active mask of threads for each side of the branch. In other words, the threads that don’t take a given branch sit idle while the active threads do their work. (Note “sit idle” might involve doing all the work and throwing away the result.)

If you have two branches, and one is trivial while the other is expensive, and if the compiler doesn’t optimize away the branch already, it may be better for performance to write the code to take both branches unconditionally, and use a conditional assignment at the end.

It’s worth knowing that often there are clever techniques to completely avoid branching. Sometimes these techniques are simple, and sometimes they’re invasive and difficult to implement. It’s easy (for me, anyway) to get stuck thinking in a single-threaded CPU way and not see how to avoid branching until you’ve bumped into and seen some of the ways smart people solve these problems.

I love seeing the codegen output, makes it easy to understand the issue, but claiming that it's faster or slower without actual benchmarks is a bit disappointing.

this.

why waste brain cells on theory when you should simply bench both versions and validate without buying into any kind of micro-optimization advice at face value.

Do shader compilers have optimization passes to undo this mistake and if not, could they be added?

I’d expect most vendors do, at least in their closed source drivers. You could also check in the mesa project if this is implemented but it’s definitely possible to do

Shader compilers tend to be very latency sensitive, so “it takes too long to run” would be a valid reason why it is not done if it is not done.

Shader compilers mostly use LLVM even though runtime is a constraint, if the pattern is common enough it’s definitely easy to match (it’s just two intrinsics after all) meaning you can do it for cheap in instcombine which you’re going to be running anyway

For some reason, I feel like this is harder to implement than you expect. The way to find out would be to get a bunch of examples of people doing this “optimizations in shader code, look at the IR generated compared to the optimal version and figure out a set of rules to detect the bad versions and transform it into a good versions. Keep in mind that in the example, the addition operators could be replaced with logical OR operators, so there are definitely multiple variations that need to be detected and corrected.

I've checked and on "certain vendors" the mix + step is actually (slightly) better: same temp usage, lower instructions/cycles.

"of course real branches happen in GPU code"

My understanding was that they don't. All executions inside a "branch" always get executed, they're simply predicated to do nothing if the condition to enter is not true.

That's only if execution is incoherent. If all threads in a warp follow the branch the same way, then all of the instructions in the not taken branch are skipped.

I've been doing this from day 1 becausee I just assumed you are not supposed to have loops or if else blocks in your shader code. Now I know better, thanks iq ur g.

What is the AMD and Microsoft cshader compiler , how do I generate and inpect these intermediate codes on my computer?

For AMD you can use the Radeon GPU Analyzer: https://gpuopen.com/rga/

> For the record, of course real branches do happen in GPU code

Well, for some definition of "real". There are hardware features (on some architectures) that implement semantics that evaluate the same way that "branched" scalar code would. There is no branching at the instruction level, and can't be on SIMD (because the other parallel shaders being evaluated by the same instructions might not have taken the same branch!)

There is real branching on the hardware level, but yes it needs to take the same branch for the whole workgroup and anything else needs to be "faked" in some form.

Yeah, exactly: I argue that a global backwards-only branch used to implement loops by checking all lanes for an "end" state is not actually a "real" branch.

It's a semantic argument, but IMHO an important one. Way, way too many users of GPUs don't understand how the code generation actually works, leading to articles like this one. Ambiguous use of terms like "branch" are the problem.

Using `mix()` isn't necessarily bad. Using boolean operations like `lessThan()` is probably better than `step()`. I just tested two ways of converting linear RGB to sRGB. On AMD, they compile to the same assembly.

Method 1.

    float linear_to_srgb(float v) {
        return v < 0.0031308 ? v * 12.92 : 1.055 * pow(v, 1.0 / 2.4) - 0.055;
    }
    
    vec3 linear_to_srgb(vec3 rgb) {
        return vec3(linear_to_srgb(rgb.r), linear_to_srgb(rgb.g), linear_to_srgb(rgb.b));
    }

Method 2:

    vec3 linear_to_srgb(vec3 rgb) {
        bvec3 cutoff = lessThan(rgb, vec3(0.0031308));
        vec3 upper = vec3(1.055) * pow(rgb, vec3(1.0 / 2.4)) - vec3(0.055);
        vec3 lower = rgb * vec3(12.92);
        return mix(upper, lower, cutoff);
    }

It is weird how long misinformation like this sticks around, the conditional move/select approach has been superior for decades on both CPU & GPU, but somehow some people still write the other approach as an "optimization".

"Conditional move is superior on CPU" is an oversimplification, in reality it was explained in 2007 by Linus[1] and nothing changed since then. Or in fact, branch predictors are constantly improving[2], while cmov data dependency problem can't be solved. Yes, cmov is better in unpredictable branches, but the universal truth is "programmers are notoriously bad at predicting how their programs actually perform"[3]

[1] https://yarchive.net/comp/linux/cmov.html

[2] https://chipsandcheese.com/p/zen-5s-2-ahead-branch-predictor...

[3] https://gcc.gnu.org/onlinedocs/gcc/Other-Builtins.html#index...

> cmov data dependency problem can't be solved

Can't is a pretty strong word.

I really only care about performance in a SIMD sense, as not using SIMD means you aren't targeting performance to begin with(GPU == SIMD, CPU SSE/AVX == SIMD)

Branches fall off vs cmov/select as your lane count increases, they are still useful but only when you are certain the probability all lanes agreeing is reasonably high.

I'm not going to second guess IQ, who is one of the greatest modern authorities on shaders, but I do have some counterarguments.

- Due to how SIMD works, it's quite likely both paths of the conditional statement get executed, so its a wash

- Most importantly, if statements look nasty on the screen. Having an if statement means a discontinuity in visuals, which means jagged and ugly pixels on the output. Of course having a step function doesnt change this, but that means the code is already in the correct form to replace it with smoothstep, which means you can interpolate between the two variations, which does look good.

Why does this keep getting downvoted? This is fundamentally true, and good advice borne of experience. At least somebody would care to weight in as to why they disagree?

I am not 100% sure nor did I downvote, but it doesn't look like a counter-argument at all.

> Due to how SIMD works, it's quite likely both paths of the conditional statement get executed, so its a wash

It's not just quite likely, it's what IQ is showing in the disassembly. For a ternary op like this one with trivial expressions on each side, the GPU evals both and then masks the result given the result of the condition.

> a step function doesnt change this, but that means the code is already in the correct form to replace it with smoothstep, which means you can interpolate between the two variations, which does look good.

A smoothstep does smooth interpolation of two values. It seems unrelated to the issue in the post. step() relates to the ternary op in the sense that both can be used to express conditionals. The post explains why you wouldn't necessarily want to use step() vs ternary op. smoothstep is related to step in some sense, but not in a way that relates to the article? i.e., going from step() to smoothstep() will entirely change the semantics of the program precisely because of the 'smooth' part.

Step vs ternary op vs if statement are equivalent - you are correct in that. What I'm saying is not about optimization or assembly - optimization doesn't matter when the end result looks bad.

What I'm saying that no matter how you express it in code, abrupt transitions of values introduce aliasing (or 'edge shimmer'), which looks unpleaseant. The way you get rid of it by smoothly blending between 2 values with smoothstep for example.

Reply on news.ycombinator.com