This argument is less effective given that SIMD is not always a straightforward substitute for vector processing. Sometimes we want 128, 256 or 512 bits of processing as a unit and will follow it up with something different, not a repeated instance of that same process.
We also used SIMD quite extensively as a 'wider GPR' - not doing stuff over tons of input characters but instead using the superior size of SIMD registers to implement things like bitwise string and NFA matchers.
A SIMD instruction can be a reasonable proxy for a wide vector processor but the reverse is not true - a specialized vector architecture is unlikely to be very helpful for this kind of 'mixed' SIMD processing. Almost any "argument from DAXPY" fails for the much richer uses of SIMD processing among active practitioners using modern SIMD.
I went on a bit of a research expedition to see if there was something that scaled better a general permutation instruction for SIMD machines. General permute scales at O(log(N)) in gate delay+ and O(N^2 * log(N)) in area, where N is the vector length. Its a full crossbar, but the fanout on the wires adds an additional log(N) factor in buffers.
For a while, it seemed like a set of instructions based on a rank-2 CLOS network (aka butterfly network) would get the job done. It scales at O(log(N)) in gate delay+ and O(N * log(N)) in area, and is very capable. Fanout is O(1). You can do all kinds of inter- and intra-lane swaps, rotations, and shifts with it. You can even do things like expand packed 16-bit RGB into 32-bit uniform RGBA.
But things like that SMH algorithm are definitely out of scope: Each input bit can only appear in at most one output location with the butterfly. So the cost to replicate a prefix scales at O(repetitions), which is unfortunate. Some algorithms based on using general shuffle are also relying on the use of PSHUFB's functionality as a complete lookup table, which the butterfly network can't do, either.
My conclusion was that you're basically stuck with a general permute instruction for a modern SIMD ISA, scaling be damned.
+ The latency scaling is somewhat misleading thanks to wire delay - they are both O(N) in wire delay.
I completely agree in principle (some applications are good fit for traditional long vector processing, but it's a terrible fit for others).
However the RISC-V Vector Extensions basically let you use the processor in "SIMD mode" by setting the vector length to a small value. It will depend on the processor architecture whether a vector length of e.g. 4 is efficient or not, but I expect for many implementations it will be relatively efficient (and you can definitely just use it as a "wider GPR" if you want to).
The only catch at the ISA level is it costs an instruction to change the vector size. So if you keep swiching between e.g. 128-bit and 512-bit instructions at a very fine granularity, that might add overhead... I'm not sure that's a very common case though?
I agree. It seems like this strategy makes only the most braindead applications of SIMD better (simple loops that can be vectorized by an arbitrary factor), but doesn't really do anything to help the meatier SIMD workloads. Most SIMD code isn't as simple as this, and the code that is usually isn't a significant factor in either developer experience or runtime.
Seriously. A lot of these proposals go veering off into second-order considerations ("Easier to decode!" "A few picojoules less energy") as I'd be very surprised if the bottlenecks are going to be from SIMD vs vector architecture ISA issues - as compared to, say, memory bandwidth or multiply-add bandwidth.
A few years ago I tried to buy a liquid cooled overclocked server for trading. Enabling AVX cost extra due to the concentrated heat output from the MMU and each core's vector unit.
It was along the lines of being able to get a server that was tested stable at 5GHz without AVX vs 4.5 GHz with AVX for the same price.
So at least on Intel, these vector units are apparently limiting clock speeds and yields due to power consumption.
Agreed, but then again, the vector architectures have the undeniable benefit of being agnostic to the underlying hardware.
With a SIMD architecture, as registers get ever wider you need to recompile your code, and binary releases require multiple code paths.
With a vector architecture, the hardware designers can just increase the machine's internal vector size and existing code will benefit immediately.
Furthermore, it seems reasonable to count GPU users among "active practitioners of modern SIMD". All code that is written in a GPU-style -- which admittedly isn't too common yet on CPUs, but it would certainly be feasible and is possible since ispc came along -- immediately becomes a beneficiary of a well-designed vector instruction architecture.
Yes, recompiling can be a pain, but there's an argument that we should get used to recompiling things from source and make sure this path is smooth, rather than trying making old binaries work as long as possible. Static binaries are another reason.
You lose binary portability though: what is you can write you image processing code such that it operates well whether your SIMD instructions is 4, 8, or 16 wide? (or more).
Then shipping this binary for mobile phones for instance would take advantage of the wider register available on high-end device while working well on more modest CPUs.
Advanced compute shaders very often need to synchronize local data within warp. When you program such shaders you have to be aware about SIMD width (32 on NVidia and Intel, 64 on AMD), and design both algorithms and data structures accordingly. Failing to do so often have significant performance costs, I saw up to 10x speedup after implementing cooperative algorithm instead of straightforward one.
I have done a bit of simd programming in past & I can tell you that its quite common to expect to rewrite your simd optimized code as a new architecture come along. sometimes you do it even as you go between different iterations of the same arch (like cortex-A8 vs A9) because of different instruction timing (and sometimes bugs). In general, asking hardware to auto-optimize around your code doesn't works except for simple problems & even then you are likely leaving performance on table.
What I really want is a lots of different types of vector & permute ops and the ability to reconfigure the simd unit on a dime when I am done with a specific type of compute (like crypto) and switch to another type (like signal processing).
Well, that makes sense if you want to squeeze the last bit of performance out of highly optimized code.
However, there's a case to be made that, in order to utilize our hardware better, we should be using vector units much more often. To make that feasible, we need a good programming paradigm that doesn't have to be rewritten for a different architecture. If that ends up not utilizing the hardware perfectly, that's okay: using a 256-bit vector unit even at 50% of the potential performance is still many times faster than scalar code.
GPU Coders haven't changed their code in the last 10 years, even as NVidia changed their architecture repeatedly.
PTX Assembly from NVidia still runs on today's architectures. I think this variable-length issue they focus on so much is a bit of a red-herring: NVidia always was 32-way SIMD but the PTX Code remains portable nonetheless.
The power is that PTX Assembly (and AMD's GCN Assembly) has a scalar-model of programming, but its execution is vectorized. So you write scalar code, but the programmer knows (and assumes it to be) in a parallel context. EDIT: I guess PTX is technically interpreted: the number of registers is not fixed, etc. etc. Nonetheless, the general "SIMD-ness" of PTX is static, and has survived a decade of hardware changes.
There are a few primitives needed for this to work: OpenCL's "Global Index" and "Local Index" for example. "Global Index" is where you are in the overall workstream, while "Local Index" is useful because intra-workgroup communications are VERY VERY FAST.
And... that's about it? Really. I guess there are a bunch of primitives (the workgroup swizzle operations, "ballot", barrier, etc. etc.), but the general GPU model is actually kinda simple.
-----------
I see a lot of these CPU architecture changes, but none of them really seem to be trying to learn from NVidia or AMD's model. A bit of PTX-assembly or GCN Assembly probably would do good to the next generation of CPU Architects.
I'm really not sure something like ptx is really the answer.
PTX isn't really native and is JIT compiled into native by the gpu driver. In many high performance apps like games native compiled versions are included for several architectures.
GPU programming model is simple because it's limited. Low single-threaded performance, high latency masked by threads switching, inefficient branching, limited memory model (can't allocate RAM, write access is very limited).
If you're happy with these limitations, write OpenCL code and run it on CPU. Will work much faster than scalar code but likely slower than a GPU would.
Low single-threaded performance, high latency masked by threads switching,
Those are fundamental design tradeoffs of the GPU architecture. Not of PTX Assembly language.
inefficient branching
Ehh? I disagree strongly. GPUs branch very efficiently, considering its within a SIMD-environment. Modern GPUs can arbitrarily diverge execution paths and reconverge as necessary. Its far more advanced than any SIMD I've seen implemented on a CPU.
GPUs have the most efficient way to branch when you're managing 32 or 64 work items at a time.
limited memory model (can't allocate RAM, write access is very limited).
You can definitely allocate RAM on a GPU. That's what the CudaMalloc function does. What do you mean by "limited write access" ??
---------
Programmers don't want to "think" in SIMD. They want to think in terms of scalar code. Per-lane Gather/Scatter is far easier to think about than vpshufb (even though vpshufb is effectively a gather instruction over an AVX register)
That's the difference. GPU Assembly language is almost equivalent to CPU stuff, its just "seen" in a different light.
CPUs are missing very, very few assembly instructions before they can run like a GPU. CPUs simply need a hardware accelerated branching-mechanism to help handle the divergent flow graphs (a simple advancement over bitmasks... AMD GCN is a good example. NVidia's Volta/Turing fully divergent flow would be nicer, but that requires a program-counter per SIMD-lane).
AMD's GCN architecture implements that mostly in just two assembly instructions: S_CBRANCH_FORK and S_CBRANCH_JOIN.
Intel is missing the "scatter-equivalent" of vpshufb. GPUs can gather/scatter arbitrarily between SIMD lanes, but Intel AVX512 only has "SIMD-gather" in the "vpshufb" instruction. Add "SIMD-scatter" and AVX512 would be complete. AMD GCN calls the two instructions "permute" and "bpermute" (backwards permute). So maybe a "backwards-vpshufb" or "vpbshufb" instruction is all AVX512 needs.
Finally, allow SIMD-instructions on CPUs to pipeline in more complex manners, to increase utilization. Only guarantee that they complete through the use of a barrier instruction. This one is an architectural shift, but GPUs are aware of "larger" workgroup sizes. (AMD GCN natively executes 64-at-a-time. But it can grow thanks to the barrier-mechanism).
The "variable length" discussed in the article here is backwards. GPUs grow their workgroups bigger, and then the programmer creates synchronization points with barrier instructions. Its actually a very simple model to work with.
Bam. Those 5 or 6 assembly instructions would make CPU-programming effectively on the same level as GPU programming.
Those are fundamental design tradeoffs of the GPU architecture. Not of PTX Assembly language.
PTX is not general purpose, it was specifically designed for that architecture, and it incorporates these tradeoffs.
far more advanced than any SIMD I've seen implemented on a CPU.
Less advanced than normal non-SIMD branching available to CPU cores. A lot of practical algorithms need both SIMD compute and scalar branching.
That's what the CudaMalloc function does.
CudaMalloc function can't be called from GPU code. You can't allocate RAM at all from inside your algorithm, there's no stack, no heap, nothing. You have to know in advance how much RAM do you need. For some practical problems this is a showstopper, e.g. try to implement unzip in CUDA.
What do you mean by "limited write access"?
You only have fixed-length arrays/textures. Global memory barriers are very slow. There're producer-consumer buffers but practically speaking they're merely thread safe cursors over statically sized buffer.
GPUs can multiply dense matrices very fast, but many practical problems are different, e.g. for sparse matrices GPUs don't deliver value performance wise. SIMD on CPU is often very useful for such problems, but yes, programming model is different, lower level and more complex. No free lunches.
Programmers don't want to "think" in SIMD.
I'm a programmer and I like SIMD.
even though vpshufb is effectively a gather instruction over an AVX register
On GPU it would be because you don't care about latency you just spawn more threads.
On CPU it's not because shuffle_epi8 is 1 cycle latency instruction, and RAM access is much slower, if you'll think they're equivalent you'll miss the performance difference.
CPUs are missing very, very few assembly instructions before they can run like a GPU
Even if you'll add these few instructions, GPUs will still be much faster, by orders of magnitude. Hardware is too different but it's not the instructions, CPU is spending transistors minimizing latency (caches and their sync, branch prediction, speculative execution, etc.) GPUs don't care about latency.
It's not instruction set that allowed simple programming model on GPUs. It's fundamentally different tradeoffs in hardware.
I think I overcomplicated my previous post. Lemme cut back the cruft and simplify. I think CPUs (Specifically AVX512) should implement the following instructions:
1. Barriers and Workgroups -- Scale SIMD UP, not downwards. The variable-length vector (discussed in this article) is backwards to the current programming model. GPU Programmers combine lanes with the concept of a OpenCL workgroup or CUDA Thread Block, and it works pretty well in my experience.
2. Implement AMD GCN-style branching with S_CBRANCH_FORK and S_CBRANCH_JOIN. This will accelerate branching when SIMD-lanes diverge in execution paths.
3. Implement "backwards vpshufb". GPUs can gather or scatter values between lanes, while CPUs can only gather data between lanes (with vpshufb). Intel AVX512 is missing an obvious and very important instruction for high-speed communication between SIMD lanes.
I agree these changes would be nice. OpenCL and similar would probably work faster on CPUs with these instructions.
I’m just not sure it’s worth it. People who are OK with GPU programming model are already using GPUs because way more powerful. AVX-512 theoretical max is 64 FLOPs/cycle, a modern $600 CPU i7-7820X with good enough cooling is capable of 1.8 TFlops single precision. A generation old $600 GPU 1080Ti is capable of 10.6 TFLops. Huge difference.
GPUs can gather or scatter values between lanes
Can they? AFAIK they can only permute values between lanes but not scatter/gather.
__shfl_sync() in CUDA does exactly the same as _mm256_permutevar8x32_ps() in AVX2 or _mm512_permutexvar_ps in AVX512.
Can they? AFAIK they can only permute values between lanes but not scatter/gather.
GPUs have a crossbar between OpenCL __shared memory and every work-item in a workgroup. Its so innate to the GPU that its almost implicit. In terms of OpenCL, the code looks like this:
The above is the equivalent to a "backwards vpshufb". AVX512 is missing this equivalent.
GPUs don't really "need" a a dedicated instruction to do this, because __local memory is a full-speed crossbar when workgroups are of size 32 (NVidia) or 64 (AMD), the native SIMD size. Its performance characteristics are not quite equivalent to vpshufb, but its still very, very, very fast in practice.
__shfl_sync() in CUDA does exactly the same as _mm256_permutevar8x32_ps() in AVX2 or _mm512_permutexvar_ps in AVX512.
There's little reason to use __shfl_sync(), because it goes through CUDA Shared memory anyway. Ditto with AMD's __amdgcn_ds_permute() and __amdgcn_ds_bpermute() intrinsics.
EDIT: I guess __shfl_sync() and the __amdgcn_ds_permute / bpermute instructions save a step. They're smaller assembly language and more concise. But I expect the overall performance to not be much different from using LDS / Shared Memory explicity.
------------------
I’m just not sure it’s worth it. People who are OK with GPU programming model are already using GPUs because way more powerful. AVX-512 theoretical max is 64 FLOPs/cycle, a modern $600 CPU i7-7820X with good enough cooling is capable of 1.8 TFlops single precision. A generation old $600 GPU 1080Ti is capable of 10.6 TFLops. Huge difference.
People don't program CPUs for FLOPs. They program on CPUs for minimum latency.
SIMD Compute is useful on a CPU because it stays in L1 cache. L1 cache is 64kB, more than enough to have a good SIMD processor accelerate some movement. CPUs even have full bandwidth to L2 cache, which is huge these days (512kB on EPYC to 1MB on Skylake-server)
CPU-based SIMD won't ever be as big or as broad as GPU-based SIMD. But... CPU-based SIMD should become easier as Intel figures out how to adopt OpenCL or CUDA programming paradigms.
There are already many problems implemented in CPU-AVX512 which execute faster than 15.8GB/s that a PCIe x16 bus will give you. Therefore, its more efficient to execute the whole problem on the CPU, rather than transfer the data to the GPU.
The data exchange is performed between registers, and more efficient than going through shared memory, which requires a load, a store and an extra register to hold the address.
If you count shared memory scatter/gather, CPU SIMD already have both. Scatter very recently so, only appeared in AVX512. Gather is available for 5 years now, _mm256_i32gather_ps was introduced in AVX2, albeit it's not particularly fast.
They program on CPUs for minimum latency.
Not just that. I code for CPU SIMD very often, and only occasionally for GPGPU. Even for code that would work very well on GPUs. The main reason for me is compatibility. I mostly work on desktop software, picking CUDA decreases userbase by a factor of 2 which is often not an option. But yeah, another reason is that CPU SIMD is fast enough already and spending time on PCIx IO doesn't pay off.
Update: another reason why I don't code GPGPU more is different programming model. GPU programming model makes writing device code easy, and like you mentioned earlier it even has good scalability built-in i.e. in many cases compute shaders need not to be aware of the warp size.
But the downside is upfront engineering costs.
I have to keep my data in very small number of continuous buffers. I have to upload these buffers to GPU. I have to know in advance how much VRAM do I need for output data.
I find this part much easier for CPU SIMD, on CPU I only need to design the lowest level of my data structures accordingly, but I can use anything at all on higher levels of the structures: hash maps, trees, linked graphs, they all work just fine, as long as their lower-level nodes are not too small, aligned, dense, and composed of these 128/256 bits SIMD vectors.
If you count shared memory scatter/gather, CPU SIMD already have both. Scatter very recently so, only appeared in AVX512. Gather is available for 5 years now, _mm256_i32gather_ps was introduced in AVX2, albeit it's not particularly fast.
You can say that again. Its still not very fast btw.
VPSCATTERDD ZMM-register is measured to be 17-clock cycles (!!!), while VPGATHERDD ZMM-register is 9-clock cycles. Gather/scatter on CPUs is very, very slow!
Its actually faster to gather/scatter through a for-loop than to actually use the VPGATHERDD or VPSCATTERDD instructions.
In contrast, Shared Memory on GPUs is a full crossbar on NVidia and AMD.
-------------
I think my details are getting a bit AMD specific. Lemme do a citation:
As a previous post briefly described, GCN3 includes two new instructions: ds_permute_b32 and ds_bpermute_b32 . They use LDS hardware to route data between the 64 lanes of a wavefront, but they don’t actually write to an LDS location.
The important tidbit is that the Load/Store units of AMD's GPU Core can support a gather/scatter to separate LDS memory banks at virtually no cost. This is a full crossbar that allows GPUs to swap lanes between SIMD registers in different lanes.
Intel CPUs do NOT have this feature, outside of vpshufb. I'm arguing that GPU Shared memory is of similar performance to vpshufb (a little bit slower, but still way faster than even a CPU's Gather/Scatter).
So yes, the "bpermute" and "permute" instructions on AMD are a full-crossbar and can execute within 2-clock cycles (if there are no bank conflicts). That's 64-dwords that can be shuffled around in just 2-clock cycles.
In contrast, Intel's Gather/scatter to L1 cache is 9-clocks or 17-clocks respectively.
-----------
The important thing here is to do a high-speed permute. The programmer can choose to use vpgatherdd, vpshufb, GPU permute, GPU bpermute, GPU LDS Memory, etc. etc. It doesn't matter for program correctness: it all does the same thing.
But GPUs have the highest-performance shuffle and permute operators. Even if you go through LDS Memory. In fact, the general permute operators of AMD GPUs just go through LDS memory, that's their underlying implementation!
Quite expected because RAM in general is much slower than registers. Even L1 cache on CPU / groupshared on GPU. On both CPUs and GPUs, you want to do as much work as possible with the data in registers.
The reason why it’s OK on GPU is massive parallelism hiding latency i.e. the hardware computes other threads instead of just waiting for data to arrive. But even on CPUs, if the requested data is not too far (in L1 or L2), hyperthreading in most modern CPUs does acceptable job in this situation.
Intel CPUs do NOT have this feature, outside of vpshufb
Yes they have. If you prefer assembly, look vpermps instruction. It can permute lanes arbitrary even across 128-bit lanes (this is unlike vpshufb/vshufps/etc.), with permutation indices being taken from another vector register. Quite fast, specifically on Haswell/Broadwell/Skylake it’s 3 cycles latency, 1 cycle throughput, single micro-op.
Stream Compaction can be used to remove redundant whitespace from strings, or to "compress" the raytracer rays so that they are all able to be read by a simple load (as opposed to a gather/scatter operation). How do you perform stream compaction?
Now, how do you do this using vpshufb? You can't. vpshufb is backwards, and won't efficiently solve this stream compaction problem. GPUs can solve this problem very efficiently, but Intel's current implementation of AVX512 is missing the "backwards" vpshufb command, to perform this operation.
Or as I've been trying to say: vpshufb is equivalent to a "gather" over SIMD Registers. But Intel is MISSING a scatter over SIMD Registers. The instruction is just... not there. I've looked for it, and it doesn't exist. As such, GPU-code (such as the stream compaction algorithm) CANNOT be implemented efficiently on a CPU.
I mean, you can use vpscatterdd to implement it, but again... vpscatterdd is 17-clock cycles. That's way too slow.
Quite expected because RAM in general is much slower than registers. Even L1 cache on CPU / groupshared on GPU. On both CPUs and GPUs, you want to do as much work as possible with the data in registers.
The L1 cache has the bandwidth to do it, it just isn't wired up correctly for this mechanism. Intel's load/store units can read/write 32-bytes at a time, in parallel across 2xload units + 1x store unit.
But the thing is: writing many small "chunks" of 4-bytes here and there (ie: a vpgatherdd / vpscatterdd operation) is not a contiguous group of 32-bytes. Therefore, Intel's cores lose a LOT of bandwidth in this case.
GPUs on the other hand, have a great-many number of load/store units. Effectively one per SIMD unit. As such, reading / writing to LDS "shared" memory on AMD GPUs can be done 32-at-a-time.
So the equivalent "vpgatherdd" over LDS cache will execute in something like 2-clock ticks on AMD GPUs (assuming no bank conflicts), while it'd take 9-clock ticks on Intel cores.
Again, LDS cache is so fast on GPUs, that it is effectively functioning as if any LDS-load/store is as fast as a vpshufb instruction. (Not quite: vpshufb is 1-clock tick, and doesn't have to worry about bank-conflicts. So vpshufb is still faster... but GPUs gather/scatter capabilities are downright incredible)
How long before you think Intel will implement a true crossbar so that the vpgatherdd and vpscatterdd instructions can actually execute quickly on a CPU?
-----------
GPUs actually implement a richer language for data-movement. Intel could very easily fix this problem by writing a "backwards vpshufb" instruction, but I'm not aware of anything that exists like that... or any plans to implement something like that.
You keep mentioning vpshufb despite it's unable to move data across 128 bit lanes.
Here's how to do that with vpermps, there's some overhead but not much, very likely much faster than RAM access even when in L1: https://stackoverflow.com/a/36951611/126995
Besides, new CPUs have AVX-512 that has vcompressps instruction just for that use case.
The L1 cache has the bandwidth to do it
It has bandwidth, but there's extra latency involved. Registers are faster.
It's the same on GPU, that's why nVidia recommends these CUDA permute lanes intrinsics over scatter/gather RAM access.
that doesn't have to be rewritten for a different architecture
CPU architectures are quite stable. SSE2 is almost 20 years old now. You can't even run modern Windows on a system which doesn't support it.
Vectorize to SSE and you'll get your 50% of potential performance. You can do it without any new paradigms, C and C++ support SSE intrinsics for decades already, other languages are catching up.
I thought it was obvious, but apparently I wasn't clear enough: what I meant with the need for rewrite is the need for rewrite to benefit from architectural changes.
If you wrote SSE2 code 20 years ago, then yes, it's still going to run. But it's not going to benefit from the advances made with AVX and AVX2.
Compare this to the GPU model, or ispc or vector instructions, where you automatically benefit from a wider machine without a rewrite of your code (and in the case of vector instructions even without a recompile).
If you want to do an analogy with GPU compute shaders, I think it is more accurate to compare to how GPUs can scale the number of cores without (potentially) the need to recompile, as long as enough blocks are scheduled.
This is orthogonal to the fact that these are warp-size aware I believe.
Sounds like you're getting fixated on the example. Short vectors would do what you want. And it's a lot easier to add specific vector instructions for particular use cases than adding N of them for each use case (for N simd sizes).
Are you getting hung up on the term 'vector'? That doesn't assume you're just doing linear algebra.
No he's not getting hung up on the term vector. He is saying that SIMD is fixed but wider width and that is not equivalent to arbitrary width. For example, SIMD shuffle type instructions have no arbitrary length equivalent.
How would you make an arbitrary shuffle between arbitrary-lengthed vectors faster than an arbitrary gather/scatter to and from L1 Cache?
Its possible to have a chip code specialized transforms or have a many-to-many crossbar when N is small (ex: 16), but when N is large (ex: 1024 elements), its no longer easy to see how to build a high-speed permute operator.
-----------
That's the thing. People only use vpshufb because its WAY faster than L1 cache. If Intel made a faster gather/scatter, there wouldn't be much point to the vpshufb instruction. But the vpshufb instruction is so fast, because its so specialized and small. It only has to worry about a 16-byte permute.
In short: we ALREADY have an arbitrary permute instruction. Its called gather/scatter. That's not what programmers want however. (I mean, programmers want a faster gather/scatter... but... vpshufb programmers use that operator only because its faster than L1 cache)
There's a front-end operating cost to the bookkeeping instructions, and a complexity cost to all the variants of the same instructions. For short vectors, the cpu can use the same hardware it uses today in SIMD, just that the SIMD work is in microcode instead of asm. The cost of fetching the permutation vector arg out of L1 isn't terribly high compared to the cost of fetch/decode on the bookkeeping instructions. And the cost of supporting all those instruction variants could be replaced with more functionality the front-end.
Seriously. Permutes get harder as you scale - VBMI on CNL is an indicator that 64-way is pretty good but it's still considerably more expensive than 4 16-way permutes on the same architecture.
There's a reason that gather is hard to do; I think if you rocked up and asked the architecture guys for a gather that was competitive with small-scale permute they would reply with the time-honored Intel putdown ("You are overpaid for whatever it is you do").
Hey, you can have 7-bit lookup tables at the byte level on AVX512VBMI (using the 2-register shuffle forms) and you can already have 6-bit lookups with 2-register 16-bit shuffles if you can play around on Skylake Server.
Mass availability of the VBMI goodies looks to be bottlenecked behind Icelake/Sunny Cove, so you'll have plenty of time to think through the implications of fast 6-bit lookup. :-)
Yeah, other examples include image codecs, such as JPEG: the DCT performed on the 8x8 blocks can benefit from SIMD, but the lanes aren't independent at all (matrix transposes, various intra-block additions).
I believe the blocks end up stored as an array-of-structs where the structs have 8*8 = 64 elements. Doing the DCT in multiple blocks requires somehow transposing this into a struct-of-arrays-like format, maybe a gather of every 64th element (likely a waste of memory bandwidth) or some sort of unpckl/unpckh-like instructions. Either way, this may impose non-trivial overhead, and so the benefits of extra parallelism are hidden.
(And, of course, that's all assuming there's enough registers, and I don't remember enough about JPEG to make a guess.)
Only if these multiple blocks fit in the 16 registers. If they won't fit and the data will be evicted to RAM, that extra parallelism will slow down the code, not speed up.
In a vector arch like the risc-v vector extension, if you want to process exactly 256 or whatever at a time, just set the configuration or vector length register and off you go with SIMD-style programming?
It is very unlikely that such a configuration will perform remotely as fast as a native SIMD implementation, unless there is some truly heroic specialization going on under the hood. Obviously the work is still possible with vector ops, in the same way that it's still possible with scalar ops too. But will it be fast? My guess is no.
It is very unlikely that such a configuration will perform remotely as fast as a native SIMD implementation
Why? AFAICS, a "real" vector ISA is mostly a superset of a SIMD ISA. What do you think is missing?
Of course, if you want to eke out the absolute maximum performance then you need to be aware of the microarchitecture you're targeting, and a length-agnostic vector ISA like SVE or RVV don't buy you that much. But I don't see how that's worse than having to redo your code whenever the vendor introduces new HW with a new SIMD ISA.
I guess one could argue that an implementation of a vector ISA targeting, say, linear algebra, would be different than an implementation focusing on maximum short-vector performance. Say, by having a vector length >> execution width, and using tricks like vector pipelining etc. to get performance rather than focusing on minimizing short-vector latency. But, that's a question of what the implementation is optimized for rather than saying what the ISA is good for, no?
Comments
This argument is less effective given that SIMD is not always a straightforward substitute for vector processing. Sometimes we want 128, 256 or 512 bits of processing as a unit and will follow it up with something different, not a repeated instance of that same process.
We had numerous different examples of this in the Hyperscan project and I broke out something similar on my blog: https://branchfree.org/2018/05/30/smh-the-swiss-army-chainsa...
We also used SIMD quite extensively as a 'wider GPR' - not doing stuff over tons of input characters but instead using the superior size of SIMD registers to implement things like bitwise string and NFA matchers.
A SIMD instruction can be a reasonable proxy for a wide vector processor but the reverse is not true - a specialized vector architecture is unlikely to be very helpful for this kind of 'mixed' SIMD processing. Almost any "argument from DAXPY" fails for the much richer uses of SIMD processing among active practitioners using modern SIMD.
I went on a bit of a research expedition to see if there was something that scaled better a general permutation instruction for SIMD machines. General permute scales at O(log(N)) in gate delay+ and O(N^2 * log(N)) in area, where N is the vector length. Its a full crossbar, but the fanout on the wires adds an additional log(N) factor in buffers.
For a while, it seemed like a set of instructions based on a rank-2 CLOS network (aka butterfly network) would get the job done. It scales at O(log(N)) in gate delay+ and O(N * log(N)) in area, and is very capable. Fanout is O(1). You can do all kinds of inter- and intra-lane swaps, rotations, and shifts with it. You can even do things like expand packed 16-bit RGB into 32-bit uniform RGBA.
But things like that SMH algorithm are definitely out of scope: Each input bit can only appear in at most one output location with the butterfly. So the cost to replicate a prefix scales at O(repetitions), which is unfortunate. Some algorithms based on using general shuffle are also relying on the use of PSHUFB's functionality as a complete lookup table, which the butterfly network can't do, either.
My conclusion was that you're basically stuck with a general permute instruction for a modern SIMD ISA, scaling be damned.
+ The latency scaling is somewhat misleading thanks to wire delay - they are both O(N) in wire delay.
I completely agree in principle (some applications are good fit for traditional long vector processing, but it's a terrible fit for others).
However the RISC-V Vector Extensions basically let you use the processor in "SIMD mode" by setting the vector length to a small value. It will depend on the processor architecture whether a vector length of e.g. 4 is efficient or not, but I expect for many implementations it will be relatively efficient (and you can definitely just use it as a "wider GPR" if you want to).
The only catch at the ISA level is it costs an instruction to change the vector size. So if you keep swiching between e.g. 128-bit and 512-bit instructions at a very fine granularity, that might add overhead... I'm not sure that's a very common case though?
I agree. It seems like this strategy makes only the most braindead applications of SIMD better (simple loops that can be vectorized by an arbitrary factor), but doesn't really do anything to help the meatier SIMD workloads. Most SIMD code isn't as simple as this, and the code that is usually isn't a significant factor in either developer experience or runtime.
Seriously. A lot of these proposals go veering off into second-order considerations ("Easier to decode!" "A few picojoules less energy") as I'd be very surprised if the bottlenecks are going to be from SIMD vs vector architecture ISA issues - as compared to, say, memory bandwidth or multiply-add bandwidth.
A few years ago I tried to buy a liquid cooled overclocked server for trading. Enabling AVX cost extra due to the concentrated heat output from the MMU and each core's vector unit.
It was along the lines of being able to get a server that was tested stable at 5GHz without AVX vs 4.5 GHz with AVX for the same price.
So at least on Intel, these vector units are apparently limiting clock speeds and yields due to power consumption.
Yes, but not due to instruction decode costs, which is really all this article is talking about.
The real heat comes from actually doing the work, not decoding what work to do.
Agreed, but then again, the vector architectures have the undeniable benefit of being agnostic to the underlying hardware.
With a SIMD architecture, as registers get ever wider you need to recompile your code, and binary releases require multiple code paths.
With a vector architecture, the hardware designers can just increase the machine's internal vector size and existing code will benefit immediately.
Furthermore, it seems reasonable to count GPU users among "active practitioners of modern SIMD". All code that is written in a GPU-style -- which admittedly isn't too common yet on CPUs, but it would certainly be feasible and is possible since ispc came along -- immediately becomes a beneficiary of a well-designed vector instruction architecture.
Yes, recompiling can be a pain, but there's an argument that we should get used to recompiling things from source and make sure this path is smooth, rather than trying making old binaries work as long as possible. Static binaries are another reason.
You lose binary portability though: what is you can write you image processing code such that it operates well whether your SIMD instructions is 4, 8, or 16 wide? (or more).
Then shipping this binary for mobile phones for instance would take advantage of the wider register available on high-end device while working well on more modest CPUs.
Shaders in games have worked this way on GPUs for probably a decade now without much issue.
Only graphics shaders.
Advanced compute shaders very often need to synchronize local data within warp. When you program such shaders you have to be aware about SIMD width (32 on NVidia and Intel, 64 on AMD), and design both algorithms and data structures accordingly. Failing to do so often have significant performance costs, I saw up to 10x speedup after implementing cooperative algorithm instead of straightforward one.
I have done a bit of simd programming in past & I can tell you that its quite common to expect to rewrite your simd optimized code as a new architecture come along. sometimes you do it even as you go between different iterations of the same arch (like cortex-A8 vs A9) because of different instruction timing (and sometimes bugs). In general, asking hardware to auto-optimize around your code doesn't works except for simple problems & even then you are likely leaving performance on table.
What I really want is a lots of different types of vector & permute ops and the ability to reconfigure the simd unit on a dime when I am done with a specific type of compute (like crypto) and switch to another type (like signal processing).
Well, that makes sense if you want to squeeze the last bit of performance out of highly optimized code.
However, there's a case to be made that, in order to utilize our hardware better, we should be using vector units much more often. To make that feasible, we need a good programming paradigm that doesn't have to be rewritten for a different architecture. If that ends up not utilizing the hardware perfectly, that's okay: using a 256-bit vector unit even at 50% of the potential performance is still many times faster than scalar code.
GPU Coders haven't changed their code in the last 10 years, even as NVidia changed their architecture repeatedly.
PTX Assembly from NVidia still runs on today's architectures. I think this variable-length issue they focus on so much is a bit of a red-herring: NVidia always was 32-way SIMD but the PTX Code remains portable nonetheless.
The power is that PTX Assembly (and AMD's GCN Assembly) has a scalar-model of programming, but its execution is vectorized. So you write scalar code, but the programmer knows (and assumes it to be) in a parallel context. EDIT: I guess PTX is technically interpreted: the number of registers is not fixed, etc. etc. Nonetheless, the general "SIMD-ness" of PTX is static, and has survived a decade of hardware changes.
There are a few primitives needed for this to work: OpenCL's "Global Index" and "Local Index" for example. "Global Index" is where you are in the overall workstream, while "Local Index" is useful because intra-workgroup communications are VERY VERY FAST.
And... that's about it? Really. I guess there are a bunch of primitives (the workgroup swizzle operations, "ballot", barrier, etc. etc.), but the general GPU model is actually kinda simple.
-----------
I see a lot of these CPU architecture changes, but none of them really seem to be trying to learn from NVidia or AMD's model. A bit of PTX-assembly or GCN Assembly probably would do good to the next generation of CPU Architects.
I'm really not sure something like ptx is really the answer. PTX isn't really native and is JIT compiled into native by the gpu driver. In many high performance apps like games native compiled versions are included for several architectures.
GPU programming model is simple because it's limited. Low single-threaded performance, high latency masked by threads switching, inefficient branching, limited memory model (can't allocate RAM, write access is very limited).
If you're happy with these limitations, write OpenCL code and run it on CPU. Will work much faster than scalar code but likely slower than a GPU would.
Those are fundamental design tradeoffs of the GPU architecture. Not of PTX Assembly language.
Ehh? I disagree strongly. GPUs branch very efficiently, considering its within a SIMD-environment. Modern GPUs can arbitrarily diverge execution paths and reconverge as necessary. Its far more advanced than any SIMD I've seen implemented on a CPU.
GPUs have the most efficient way to branch when you're managing 32 or 64 work items at a time.
You can definitely allocate RAM on a GPU. That's what the CudaMalloc function does. What do you mean by "limited write access" ??
---------
Programmers don't want to "think" in SIMD. They want to think in terms of scalar code. Per-lane Gather/Scatter is far easier to think about than vpshufb (even though vpshufb is effectively a gather instruction over an AVX register)
That's the difference. GPU Assembly language is almost equivalent to CPU stuff, its just "seen" in a different light.
CPUs are missing very, very few assembly instructions before they can run like a GPU. CPUs simply need a hardware accelerated branching-mechanism to help handle the divergent flow graphs (a simple advancement over bitmasks... AMD GCN is a good example. NVidia's Volta/Turing fully divergent flow would be nicer, but that requires a program-counter per SIMD-lane).
AMD's GCN architecture implements that mostly in just two assembly instructions: S_CBRANCH_FORK and S_CBRANCH_JOIN.
Intel is missing the "scatter-equivalent" of vpshufb. GPUs can gather/scatter arbitrarily between SIMD lanes, but Intel AVX512 only has "SIMD-gather" in the "vpshufb" instruction. Add "SIMD-scatter" and AVX512 would be complete. AMD GCN calls the two instructions "permute" and "bpermute" (backwards permute). So maybe a "backwards-vpshufb" or "vpbshufb" instruction is all AVX512 needs.
Finally, allow SIMD-instructions on CPUs to pipeline in more complex manners, to increase utilization. Only guarantee that they complete through the use of a barrier instruction. This one is an architectural shift, but GPUs are aware of "larger" workgroup sizes. (AMD GCN natively executes 64-at-a-time. But it can grow thanks to the barrier-mechanism).
The "variable length" discussed in the article here is backwards. GPUs grow their workgroups bigger, and then the programmer creates synchronization points with barrier instructions. Its actually a very simple model to work with.
Bam. Those 5 or 6 assembly instructions would make CPU-programming effectively on the same level as GPU programming.
PTX is not general purpose, it was specifically designed for that architecture, and it incorporates these tradeoffs.
Less advanced than normal non-SIMD branching available to CPU cores. A lot of practical algorithms need both SIMD compute and scalar branching.
CudaMalloc function can't be called from GPU code. You can't allocate RAM at all from inside your algorithm, there's no stack, no heap, nothing. You have to know in advance how much RAM do you need. For some practical problems this is a showstopper, e.g. try to implement unzip in CUDA.
You only have fixed-length arrays/textures. Global memory barriers are very slow. There're producer-consumer buffers but practically speaking they're merely thread safe cursors over statically sized buffer.
GPUs can multiply dense matrices very fast, but many practical problems are different, e.g. for sparse matrices GPUs don't deliver value performance wise. SIMD on CPU is often very useful for such problems, but yes, programming model is different, lower level and more complex. No free lunches.
I'm a programmer and I like SIMD.
On GPU it would be because you don't care about latency you just spawn more threads.
On CPU it's not because shuffle_epi8 is 1 cycle latency instruction, and RAM access is much slower, if you'll think they're equivalent you'll miss the performance difference.
Even if you'll add these few instructions, GPUs will still be much faster, by orders of magnitude. Hardware is too different but it's not the instructions, CPU is spending transistors minimizing latency (caches and their sync, branch prediction, speculative execution, etc.) GPUs don't care about latency.
It's not instruction set that allowed simple programming model on GPUs. It's fundamentally different tradeoffs in hardware.
I think I overcomplicated my previous post. Lemme cut back the cruft and simplify. I think CPUs (Specifically AVX512) should implement the following instructions:
1. Barriers and Workgroups -- Scale SIMD UP, not downwards. The variable-length vector (discussed in this article) is backwards to the current programming model. GPU Programmers combine lanes with the concept of a OpenCL workgroup or CUDA Thread Block, and it works pretty well in my experience.
2. Implement AMD GCN-style branching with S_CBRANCH_FORK and S_CBRANCH_JOIN. This will accelerate branching when SIMD-lanes diverge in execution paths.
3. Implement "backwards vpshufb". GPUs can gather or scatter values between lanes, while CPUs can only gather data between lanes (with vpshufb). Intel AVX512 is missing an obvious and very important instruction for high-speed communication between SIMD lanes.
I agree these changes would be nice. OpenCL and similar would probably work faster on CPUs with these instructions.
I’m just not sure it’s worth it. People who are OK with GPU programming model are already using GPUs because way more powerful. AVX-512 theoretical max is 64 FLOPs/cycle, a modern $600 CPU i7-7820X with good enough cooling is capable of 1.8 TFlops single precision. A generation old $600 GPU 1080Ti is capable of 10.6 TFLops. Huge difference.
Can they? AFAIK they can only permute values between lanes but not scatter/gather.
__shfl_sync() in CUDA does exactly the same as _mm256_permutevar8x32_ps() in AVX2 or _mm512_permutexvar_ps in AVX512.
GPUs have a crossbar between OpenCL __shared memory and every work-item in a workgroup. Its so innate to the GPU that its almost implicit. In terms of OpenCL, the code looks like this:
The above is roughly equivalent to vpshufb, where generate_index() is the parameter to vpshufb. The above is the equivalent to a "backwards vpshufb". AVX512 is missing this equivalent.GPUs don't really "need" a a dedicated instruction to do this, because __local memory is a full-speed crossbar when workgroups are of size 32 (NVidia) or 64 (AMD), the native SIMD size. Its performance characteristics are not quite equivalent to vpshufb, but its still very, very, very fast in practice.
There's little reason to use __shfl_sync(), because it goes through CUDA Shared memory anyway. Ditto with AMD's __amdgcn_ds_permute() and __amdgcn_ds_bpermute() intrinsics.
EDIT: I guess __shfl_sync() and the __amdgcn_ds_permute / bpermute instructions save a step. They're smaller assembly language and more concise. But I expect the overall performance to not be much different from using LDS / Shared Memory explicity.
------------------
People don't program CPUs for FLOPs. They program on CPUs for minimum latency.
SIMD Compute is useful on a CPU because it stays in L1 cache. L1 cache is 64kB, more than enough to have a good SIMD processor accelerate some movement. CPUs even have full bandwidth to L2 cache, which is huge these days (512kB on EPYC to 1MB on Skylake-server)
CPU-based SIMD won't ever be as big or as broad as GPU-based SIMD. But... CPU-based SIMD should become easier as Intel figures out how to adopt OpenCL or CUDA programming paradigms.
There are already many problems implemented in CPU-AVX512 which execute faster than 15.8GB/s that a PCIe x16 bus will give you. Therefore, its more efficient to execute the whole problem on the CPU, rather than transfer the data to the GPU.
NVidia says the opposite is true. Here's a link: https://devblogs.nvidia.com/using-cuda-warp-level-primitives...
The data exchange is performed between registers, and more efficient than going through shared memory, which requires a load, a store and an extra register to hold the address.
If you count shared memory scatter/gather, CPU SIMD already have both. Scatter very recently so, only appeared in AVX512. Gather is available for 5 years now, _mm256_i32gather_ps was introduced in AVX2, albeit it's not particularly fast.
Not just that. I code for CPU SIMD very often, and only occasionally for GPGPU. Even for code that would work very well on GPUs. The main reason for me is compatibility. I mostly work on desktop software, picking CUDA decreases userbase by a factor of 2 which is often not an option. But yeah, another reason is that CPU SIMD is fast enough already and spending time on PCIx IO doesn't pay off.
Update: another reason why I don't code GPGPU more is different programming model. GPU programming model makes writing device code easy, and like you mentioned earlier it even has good scalability built-in i.e. in many cases compute shaders need not to be aware of the warp size.
But the downside is upfront engineering costs.
I have to keep my data in very small number of continuous buffers. I have to upload these buffers to GPU. I have to know in advance how much VRAM do I need for output data.
I find this part much easier for CPU SIMD, on CPU I only need to design the lowest level of my data structures accordingly, but I can use anything at all on higher levels of the structures: hash maps, trees, linked graphs, they all work just fine, as long as their lower-level nodes are not too small, aligned, dense, and composed of these 128/256 bits SIMD vectors.
You can say that again. Its still not very fast btw.
https://www.agner.org/optimize/instruction_tables.pdf
VPSCATTERDD ZMM-register is measured to be 17-clock cycles (!!!), while VPGATHERDD ZMM-register is 9-clock cycles. Gather/scatter on CPUs is very, very slow!
Its actually faster to gather/scatter through a for-loop than to actually use the VPGATHERDD or VPSCATTERDD instructions.
In contrast, Shared Memory on GPUs is a full crossbar on NVidia and AMD.
-------------
I think my details are getting a bit AMD specific. Lemme do a citation:
https://gpuopen.com/amd-gcn-assembly-cross-lane-operations/
The important tidbit is that the Load/Store units of AMD's GPU Core can support a gather/scatter to separate LDS memory banks at virtually no cost. This is a full crossbar that allows GPUs to swap lanes between SIMD registers in different lanes.
Intel CPUs do NOT have this feature, outside of vpshufb. I'm arguing that GPU Shared memory is of similar performance to vpshufb (a little bit slower, but still way faster than even a CPU's Gather/Scatter).
So yes, the "bpermute" and "permute" instructions on AMD are a full-crossbar and can execute within 2-clock cycles (if there are no bank conflicts). That's 64-dwords that can be shuffled around in just 2-clock cycles.
In contrast, Intel's Gather/scatter to L1 cache is 9-clocks or 17-clocks respectively.
-----------
The important thing here is to do a high-speed permute. The programmer can choose to use vpgatherdd, vpshufb, GPU permute, GPU bpermute, GPU LDS Memory, etc. etc. It doesn't matter for program correctness: it all does the same thing.
But GPUs have the highest-performance shuffle and permute operators. Even if you go through LDS Memory. In fact, the general permute operators of AMD GPUs just go through LDS memory, that's their underlying implementation!
Quite expected because RAM in general is much slower than registers. Even L1 cache on CPU / groupshared on GPU. On both CPUs and GPUs, you want to do as much work as possible with the data in registers.
The reason why it’s OK on GPU is massive parallelism hiding latency i.e. the hardware computes other threads instead of just waiting for data to arrive. But even on CPUs, if the requested data is not too far (in L1 or L2), hyperthreading in most modern CPUs does acceptable job in this situation.
Yes they have. If you prefer assembly, look vpermps instruction. It can permute lanes arbitrary even across 128-bit lanes (this is unlike vpshufb/vshufps/etc.), with permutation indices being taken from another vector register. Quite fast, specifically on Haswell/Broadwell/Skylake it’s 3 cycles latency, 1 cycle throughput, single micro-op.
No they don't, but its very tricky to see why. You're blinded by vpshufb, and can't see how it can fail to solve some problems.
Lets take stream compaction as an example problem.
http://www.cse.chalmers.se/%7Euffe/streamcompaction.pdf
Stream Compaction can be used to remove redundant whitespace from strings, or to "compress" the raytracer rays so that they are all able to be read by a simple load (as opposed to a gather/scatter operation). How do you perform stream compaction?
Well, its a straightforward scatter operation: https://i.imgur.com/aIoO8dm.png
Now, how do you do this using vpshufb? You can't. vpshufb is backwards, and won't efficiently solve this stream compaction problem. GPUs can solve this problem very efficiently, but Intel's current implementation of AVX512 is missing the "backwards" vpshufb command, to perform this operation.
Or as I've been trying to say: vpshufb is equivalent to a "gather" over SIMD Registers. But Intel is MISSING a scatter over SIMD Registers. The instruction is just... not there. I've looked for it, and it doesn't exist. As such, GPU-code (such as the stream compaction algorithm) CANNOT be implemented efficiently on a CPU.
I mean, you can use vpscatterdd to implement it, but again... vpscatterdd is 17-clock cycles. That's way too slow.
The L1 cache has the bandwidth to do it, it just isn't wired up correctly for this mechanism. Intel's load/store units can read/write 32-bytes at a time, in parallel across 2xload units + 1x store unit.
But the thing is: writing many small "chunks" of 4-bytes here and there (ie: a vpgatherdd / vpscatterdd operation) is not a contiguous group of 32-bytes. Therefore, Intel's cores lose a LOT of bandwidth in this case.
GPUs on the other hand, have a great-many number of load/store units. Effectively one per SIMD unit. As such, reading / writing to LDS "shared" memory on AMD GPUs can be done 32-at-a-time.
So the equivalent "vpgatherdd" over LDS cache will execute in something like 2-clock ticks on AMD GPUs (assuming no bank conflicts), while it'd take 9-clock ticks on Intel cores.
Again, LDS cache is so fast on GPUs, that it is effectively functioning as if any LDS-load/store is as fast as a vpshufb instruction. (Not quite: vpshufb is 1-clock tick, and doesn't have to worry about bank-conflicts. So vpshufb is still faster... but GPUs gather/scatter capabilities are downright incredible)
How long before you think Intel will implement a true crossbar so that the vpgatherdd and vpscatterdd instructions can actually execute quickly on a CPU?
-----------
GPUs actually implement a richer language for data-movement. Intel could very easily fix this problem by writing a "backwards vpshufb" instruction, but I'm not aware of anything that exists like that... or any plans to implement something like that.
You keep mentioning vpshufb despite it's unable to move data across 128 bit lanes.
Here's how to do that with vpermps, there's some overhead but not much, very likely much faster than RAM access even when in L1: https://stackoverflow.com/a/36951611/126995
Besides, new CPUs have AVX-512 that has vcompressps instruction just for that use case.
It has bandwidth, but there's extra latency involved. Registers are faster.
It's the same on GPU, that's why nVidia recommends these CUDA permute lanes intrinsics over scatter/gather RAM access.
Here's a recent article about people actually using these permute intrinsics achieving quite good results: https://news.ycombinator.com/item?id=19018240
Hmm, I'll have to study AVX512 more then.
Thanks for the discussion!
You’re welcome. I’ve also learned something new. I program CPU SIMD and graphics shaders a lot, but compute shaders or CUDA only occasionally.
Ptx is NOT architecture specific and is jit into native by gpu driver...
CPU architectures are quite stable. SSE2 is almost 20 years old now. You can't even run modern Windows on a system which doesn't support it.
Vectorize to SSE and you'll get your 50% of potential performance. You can do it without any new paradigms, C and C++ support SSE intrinsics for decades already, other languages are catching up.
I thought it was obvious, but apparently I wasn't clear enough: what I meant with the need for rewrite is the need for rewrite to benefit from architectural changes.
If you wrote SSE2 code 20 years ago, then yes, it's still going to run. But it's not going to benefit from the advances made with AVX and AVX2.
Compare this to the GPU model, or ispc or vector instructions, where you automatically benefit from a wider machine without a rewrite of your code (and in the case of vector instructions even without a recompile).
If you want to do an analogy with GPU compute shaders, I think it is more accurate to compare to how GPUs can scale the number of cores without (potentially) the need to recompile, as long as enough blocks are scheduled.
This is orthogonal to the fact that these are warp-size aware I believe.
Sounds like you're getting fixated on the example. Short vectors would do what you want. And it's a lot easier to add specific vector instructions for particular use cases than adding N of them for each use case (for N simd sizes).
Are you getting hung up on the term 'vector'? That doesn't assume you're just doing linear algebra.
No he's not getting hung up on the term vector. He is saying that SIMD is fixed but wider width and that is not equivalent to arbitrary width. For example, SIMD shuffle type instructions have no arbitrary length equivalent.
Why not? You take the new ordering as an input vector.
How would you make an arbitrary shuffle between arbitrary-lengthed vectors faster than an arbitrary gather/scatter to and from L1 Cache?
Its possible to have a chip code specialized transforms or have a many-to-many crossbar when N is small (ex: 16), but when N is large (ex: 1024 elements), its no longer easy to see how to build a high-speed permute operator.
-----------
That's the thing. People only use vpshufb because its WAY faster than L1 cache. If Intel made a faster gather/scatter, there wouldn't be much point to the vpshufb instruction. But the vpshufb instruction is so fast, because its so specialized and small. It only has to worry about a 16-byte permute.
In short: we ALREADY have an arbitrary permute instruction. Its called gather/scatter. That's not what programmers want however. (I mean, programmers want a faster gather/scatter... but... vpshufb programmers use that operator only because its faster than L1 cache)
Specialize for the case of small vectors.
There's a front-end operating cost to the bookkeeping instructions, and a complexity cost to all the variants of the same instructions. For short vectors, the cpu can use the same hardware it uses today in SIMD, just that the SIMD work is in microcode instead of asm. The cost of fetching the permutation vector arg out of L1 isn't terribly high compared to the cost of fetch/decode on the bookkeeping instructions. And the cost of supporting all those instruction variants could be replaced with more functionality the front-end.
Seriously. Permutes get harder as you scale - VBMI on CNL is an indicator that 64-way is pretty good but it's still considerably more expensive than 4 16-way permutes on the same architecture.
There's a reason that gather is hard to do; I think if you rocked up and asked the architecture guys for a gather that was competitive with small-scale permute they would reply with the time-honored Intel putdown ("You are overpaid for whatever it is you do").
And now I'm distracted looking for 6-bit lookup tables that will enjoy that instruction. DES had 6-bit SBoxes for example.
https://en.wikipedia.org/wiki/DES_supplementary_material#Sub...
Hmmmmm... 6-bit lookup tables. Yum. I wonder what else is out there that would benefit?
Hey, you can have 7-bit lookup tables at the byte level on AVX512VBMI (using the 2-register shuffle forms) and you can already have 6-bit lookups with 2-register 16-bit shuffles if you can play around on Skylake Server.
Mass availability of the VBMI goodies looks to be bottlenecked behind Icelake/Sunny Cove, so you'll have plenty of time to think through the implications of fast 6-bit lookup. :-)
Yeah, other examples include image codecs, such as JPEG: the DCT performed on the 8x8 blocks can benefit from SIMD, but the lanes aren't independent at all (matrix transposes, various intra-block additions).
Do it across blocks and you can squeeze out more parallelism.
I believe the blocks end up stored as an array-of-structs where the structs have 8*8 = 64 elements. Doing the DCT in multiple blocks requires somehow transposing this into a struct-of-arrays-like format, maybe a gather of every 64th element (likely a waste of memory bandwidth) or some sort of unpckl/unpckh-like instructions. Either way, this may impose non-trivial overhead, and so the benefits of extra parallelism are hidden.
(And, of course, that's all assuming there's enough registers, and I don't remember enough about JPEG to make a guess.)
Only if these multiple blocks fit in the 16 registers. If they won't fit and the data will be evicted to RAM, that extra parallelism will slow down the code, not speed up.
In a vector arch like the risc-v vector extension, if you want to process exactly 256 or whatever at a time, just set the configuration or vector length register and off you go with SIMD-style programming?
It is very unlikely that such a configuration will perform remotely as fast as a native SIMD implementation, unless there is some truly heroic specialization going on under the hood. Obviously the work is still possible with vector ops, in the same way that it's still possible with scalar ops too. But will it be fast? My guess is no.
Why? AFAICS, a "real" vector ISA is mostly a superset of a SIMD ISA. What do you think is missing?
Of course, if you want to eke out the absolute maximum performance then you need to be aware of the microarchitecture you're targeting, and a length-agnostic vector ISA like SVE or RVV don't buy you that much. But I don't see how that's worse than having to redo your code whenever the vendor introduces new HW with a new SIMD ISA.
I guess one could argue that an implementation of a vector ISA targeting, say, linear algebra, would be different than an implementation focusing on maximum short-vector performance. Say, by having a vector length >> execution width, and using tricks like vector pipelining etc. to get performance rather than focusing on minimizing short-vector latency. But, that's a question of what the implementation is optimized for rather than saying what the ISA is good for, no?
Agreed. SIMDs true competitor is ironically the super-scalar architecture itself and its relative the VLIW.