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