3 ms·

I wish someone would make an explanation of the current NVIDIA and AMD architectures with the terms used translated to their CPU equivalents, instead of the cur

by gcp 15y ago

I wish someone would make an explanation of the current NVIDIA and AMD architectures with the terms used translated to their CPU equivalents, instead of the current half-marketing-speak which uses similar terms for different things between the two vendors. I'm finding it exceedingly hard to make any guess whether certain algorithms are worth trying to port over or not, because the explanations are almost incomprehensible. To give an example, the VLIW->SIMD unit transition in this architecture compared to NVIDIA's scalar units. As far as I understand the VLIW vs SIMD difference, a VLIW instruction is more powerful compared to SIMD because one "long instruction" can contain different operations on different data. Whereas SIMD is the same operation on multiple data. Traditionally, VLIW was entirely statically scheduled, which put all the burden on the compiler. Because graphics cards recompile all shaders anyway, it's not a bad fit. So, now AMD changed their units from 16x4 VLIW to groups of 16xSIMD engines. The advantage here is that you no longer have to have groups of 64 similar ALU operations, but can do with groups of 16 similar ALU operations. Conversely, there should be more of such groups, i.e. more control logic compared to the old design. To top that off, there's improvements that allow one to schedule multiple threads over the GPU at once. Am I on the right track here? If I am, whats the minimum amount of identical operations that you must do to achieve reasonable throughput? 16 identical operations for 100% efficiency? Put differently, if I have code that requires different operations (due to branches, i.e. effectively conditional operations) on each computation stream that I'm pulling through, what factors are going to limit the effective speed I get out of this?

DiabloD3 15y ago
VLIW merely means instructions are setup in blocks with clearly defined starts and ends (AMD calls these blocks 'clauses'). This technique is also called MIMD (multiple instruction multiple data, as opposed to SIMD, single instruction multiple data). VLIW exists to exploit very wide instruction parallelism while clearly delineating what instructions have dependencies on other instructions, and (at least on the Radeon) usually end on a memory write or a complex branch. Radeon clauses are up to 128 instructions long, and manage 5 ALUs (on 4xxx, 5xxx, and 68xx, 4 of them identical, 1 of them being able to do double precision math and transcendentals), or 4 ALUs (on 69xx, all 4 identical and does not require a specialized ALU for DP and T). The compiler optimizes dependency flow and ALU usage, instead of requiring dedicated hardware common in CPU design. This means far less silicon is dedicated to the task, and instruction scheduling is far more predictable and optimized. Your suggestion that AMD used 16 ALUs per pipe under VLIW is wrong, The change is 4/5 VLIW ALUs to 16 ALUs that now can execute SIMD instructions. The compute units (the head end that synchronizes multiple pipes to perform one task in parallel) still use VLIW-like clauses to synchronize pipe usage. Your suggestion that this new arch allows you to schedule multiple threads on the GPU at once is nonsensical: the correct term for pipe is "hardware thread", on a GPU like 5870 you have 320 hardware threads (1600 ALUs), you already schedule all of them at the same time for massively parallel execution. What has changed is the CUs now support running clauses from different shaders at the same time by using some CU on one task, and some on another, and I believe it may also be able to have clauses from different shaders loaded at the same time and switch without overhead; on VLIW Radeons, shader change out has a high context switch penalty. The only thing the new GCN arch really does is allow the ALUs to operate on SIMD instructions which allows higher instruction packing. This does not mean they do not use VLIW-like clauses, and it doesn't mean it is like Nvidia's design (which the media keeps repeating). Nvidia's ALUs are free form stream processors, and do not have a clear beginning or end to each clause (as the hardware does not exploit hardware thread synchronization), they frequently suffer from cache misses and pipeline stalls, and they cannot easily exploit instruction level parallelism. In addition, Nvidia does not exploit deep pipelining. On VLIW Radeons, the instruction pipeline is multistaged and 4 instructions deep, so by the time you are submitting the 5th instruction you are getting the results from the first and instructions 2, 3, and 4 are still being processed. This allows much easier synchronization between ALUs since they all read/write to the same set of registers. The addition of SIMD instructions to this design allows much higher data throughput and much higher instruction packing; instead of executing, say, two Bitcoin hashes per VLIW4/5 group, each instruction winding around the group for maximum ALU efficiency, you can run 4 (or however many GCN uses for SIMD, most likely 4 or 8) hashes at the same time as a SIMD operation and not require complex compiler maneuvering to do a clearly instruction parallel operation (thus 16x4 hashes per group). Now, ultimately, nothing of what I've written actually matters. OpenCL is a black box on purpose, it doesn't matter how the implementation executes it as long as it does so correctly and efficiently. AMD is betting that GCN is more efficient for the given silicon real estate.
- Tuna-Fish 15y ago
  > Your suggestion that AMD used 16 ALUs per pipe under VLIW is wrong, The change is 4/5 VLIW ALUs to 16 ALUs that now can execute SIMD instructions. Actually, the previous architecture was a setup of 16x simd, where each of the simd operations was a 5/4 wide vliw. So calling that 320 hardware threads is wrong -- in Cypress there really was only 20 front-ends which drove these bundles of 80 alus in groups of 5x16. Also, it was a 4-long barrel processor, so you had to schedule a SIMD "wavefront" of 64 "threads" for each unit. In the new version each CU still has 4x16 ALUS like it had in Cayman, but now each of the 4 simd units of 16 elements can be scheduled independently by a different hardware thread.
  - DiabloD3 15y ago
    There are 20 CUs, but it controls multiple sets of VLIW5/4 in parallel. AMD claims 1600 ALUs on the 5870, so 1600/5 = 320. I'm not sure how they are are factoring in the 4 deep pipeline for the barrel, but I'm pretty sure it doesn't mean there is only 80 actual ALUs. The R700 Programming Manual seems to indicate my interpretation is correct, although if you can provide evidence that I'm misinterpreting it, I'm all ears.
    - Tuna-Fish 15y ago
      There are 1600 Alus, but they are grouped in vliws of 5 elements, which are grouped in simd groups of 16 vliws. So there are 20 front ends, and reading a single bundle in one of them will instantly make 80 alus execute an instruction. The barrel is essentially used to extend the vector registers from 16-elem to 64-elem, and a 64 "thread" wavefront, consisting of 5 VLIW'd instructions is essentially the smallest amount of work that R700 can do.
gmt2027 15y ago
My experience is in porting numerical algorithms to NVIDIA GPUs with the CUDA platform. If your intention is to decide whether a GPU port is worth it, then perhaps you should be thinking at a higher level of abstraction. I only "speak" CUDA, however the primitives provided by the CUDA/OpenCL development platforms are largely equivalent. For a single "kernel" the GPU provides parallelism at two levels, 1. Fine-grained Thread-level parallelism: You have a group of threads performing ideally identical operations. The model is called SIMT (Single Instruction Multiple Thread) because it is not quite SIMD but close. Say you have (1-1024) threads, the hardware splits these into warps of 32 threads. The 32 threads within a warp execute together in classic SIMD manner, however only one of these warps is executing on the actual hardware at any given time. When a warp makes a request for data from high-latency device memory, it is typically suspended and another warp swapped in. This allows latency hiding between warps. However, calling __syncthreads() forces all warps to synchronise. Any conditional statements cause warp divergence, the warp is split and the alternative execution paths run sequentially. Clearly, this incurs a performance penalty as SIMD behaviour is lost and should be avoided as much as possible. The 'ideal' identical operation size is therefore a warp or 32 consecutive threads, however occasional divergence is acceptable. Memory access patterns are much more important however. 2. Coarse-grained Block-level parallelism. A block is an independent group of threads. The platform allows mutiple blocks to execute on different Streaming Multiprocessors. The hardware may schedule multiple blocks to execute on a single SM if there are sufficient resources (such as register and shared memory), this also facilitates latency hiding. Other than the kernel invocation which created multiple threads blocks, they are unrelated and cannot communicate or be synchronised. The potential thread level parallelism is what should be considered. If your algorithm is representable as fine-grained floating point vector operations without excessive divergence then it is well-suited to the GPU. Of course, the actual performance gain depends on the quality of implementation and optimisation is highly GPU-specific and non-trivial. I recommend http://developer.download.nvidia.com/compute/DevZone/C/html/C/src/reduction/doc/reduction.pdf http://developer.download.nvidia.com/compute/DevZone/C/html/... as an example of how important optimisations lead to those impressive performance numbers.

by gcp 15y ago

DiabloD3 15y ago

VLIW merely means instructions are setup in blocks with clearly defined starts and ends (AMD calls these blocks 'clauses'). This technique is also called MIMD (multiple instruction multiple data, as opposed to SIMD, single instruction multiple data). VLIW exists to exploit very wide instruction parallelism while clearly delineating what instructions have dependencies on other instructions, and (at least on the Radeon) usually end on a memory write or a complex branch. Radeon clauses are up to 128 instructions long, and manage 5 ALUs (on 4xxx, 5xxx, and 68xx, 4 of them identical, 1 of them being able to do double precision math and transcendentals), or 4 ALUs (on 69xx, all 4 identical and does not require a specialized ALU for DP and T). The compiler optimizes dependency flow and ALU usage, instead of requiring dedicated hardware common in CPU design. This means far less silicon is dedicated to the task, and instruction scheduling is far more predictable and optimized. Your suggestion that AMD used 16 ALUs per pipe under VLIW is wrong, The change is 4/5 VLIW ALUs to 16 ALUs that now can execute SIMD instructions. The compute units (the head end that synchronizes multiple pipes to perform one task in parallel) still use VLIW-like clauses to synchronize pipe usage. Your suggestion that this new arch allows you to schedule multiple threads on the GPU at once is nonsensical: the correct term for pipe is "hardware thread", on a GPU like 5870 you have 320 hardware threads (1600 ALUs), you already schedule all of them at the same time for massively parallel execution. What has changed is the CUs now support running clauses from different shaders at the same time by using some CU on one task, and some on another, and I believe it may also be able to have clauses from different shaders loaded at the same time and switch without overhead; on VLIW Radeons, shader change out has a high context switch penalty. The only thing the new GCN arch really does is allow the ALUs to operate on SIMD instructions which allows higher instruction packing. This does not mean they do not use VLIW-like clauses, and it doesn't mean it is like Nvidia's design (which the media keeps repeating). Nvidia's ALUs are free form stream processors, and do not have a clear beginning or end to each clause (as the hardware does not exploit hardware thread synchronization), they frequently suffer from cache misses and pipeline stalls, and they cannot easily exploit instruction level parallelism. In addition, Nvidia does not exploit deep pipelining. On VLIW Radeons, the instruction pipeline is multistaged and 4 instructions deep, so by the time you are submitting the 5th instruction you are getting the results from the first and instructions 2, 3, and 4 are still being processed. This allows much easier synchronization between ALUs since they all read/write to the same set of registers. The addition of SIMD instructions to this design allows much higher data throughput and much higher instruction packing; instead of executing, say, two Bitcoin hashes per VLIW4/5 group, each instruction winding around the group for maximum ALU efficiency, you can run 4 (or however many GCN uses for SIMD, most likely 4 or 8) hashes at the same time as a SIMD operation and not require complex compiler maneuvering to do a clearly instruction parallel operation (thus 16x4 hashes per group). Now, ultimately, nothing of what I've written actually matters. OpenCL is a black box on purpose, it doesn't matter how the implementation executes it as long as it does so correctly and efficiently. AMD is betting that GCN is more efficient for the given silicon real estate.

Tuna-Fish 15y ago
> Your suggestion that AMD used 16 ALUs per pipe under VLIW is wrong, The change is 4/5 VLIW ALUs to 16 ALUs that now can execute SIMD instructions. Actually, the previous architecture was a setup of 16x simd, where each of the simd operations was a 5/4 wide vliw. So calling that 320 hardware threads is wrong -- in Cypress there really was only 20 front-ends which drove these bundles of 80 alus in groups of 5x16. Also, it was a 4-long barrel processor, so you had to schedule a SIMD "wavefront" of 64 "threads" for each unit. In the new version each CU still has 4x16 ALUS like it had in Cayman, but now each of the 4 simd units of 16 elements can be scheduled independently by a different hardware thread.
- DiabloD3 15y ago
  There are 20 CUs, but it controls multiple sets of VLIW5/4 in parallel. AMD claims 1600 ALUs on the 5870, so 1600/5 = 320. I'm not sure how they are are factoring in the 4 deep pipeline for the barrel, but I'm pretty sure it doesn't mean there is only 80 actual ALUs. The R700 Programming Manual seems to indicate my interpretation is correct, although if you can provide evidence that I'm misinterpreting it, I'm all ears.
  - Tuna-Fish 15y ago
    There are 1600 Alus, but they are grouped in vliws of 5 elements, which are grouped in simd groups of 16 vliws. So there are 20 front ends, and reading a single bundle in one of them will instantly make 80 alus execute an instruction. The barrel is essentially used to extend the vector registers from 16-elem to 64-elem, and a 64 "thread" wavefront, consisting of 5 VLIW'd instructions is essentially the smallest amount of work that R700 can do.

gmt2027 15y ago

My experience is in porting numerical algorithms to NVIDIA GPUs with the CUDA platform. If your intention is to decide whether a GPU port is worth it, then perhaps you should be thinking at a higher level of abstraction. I only "speak" CUDA, however the primitives provided by the CUDA/OpenCL development platforms are largely equivalent. For a single "kernel" the GPU provides parallelism at two levels, 1. Fine-grained Thread-level parallelism: You have a group of threads performing ideally identical operations. The model is called SIMT (Single Instruction Multiple Thread) because it is not quite SIMD but close. Say you have (1-1024) threads, the hardware splits these into warps of 32 threads. The 32 threads within a warp execute together in classic SIMD manner, however only one of these warps is executing on the actual hardware at any given time. When a warp makes a request for data from high-latency device memory, it is typically suspended and another warp swapped in. This allows latency hiding between warps. However, calling __syncthreads() forces all warps to synchronise. Any conditional statements cause warp divergence, the warp is split and the alternative execution paths run sequentially. Clearly, this incurs a performance penalty as SIMD behaviour is lost and should be avoided as much as possible. The 'ideal' identical operation size is therefore a warp or 32 consecutive threads, however occasional divergence is acceptable. Memory access patterns are much more important however. 2. Coarse-grained Block-level parallelism. A block is an independent group of threads. The platform allows mutiple blocks to execute on different Streaming Multiprocessors. The hardware may schedule multiple blocks to execute on a single SM if there are sufficient resources (such as register and shared memory), this also facilitates latency hiding. Other than the kernel invocation which created multiple threads blocks, they are unrelated and cannot communicate or be synchronised. The potential thread level parallelism is what should be considered. If your algorithm is representable as fine-grained floating point vector operations without excessive divergence then it is well-suited to the GPU. Of course, the actual performance gain depends on the quality of implementation and optimisation is highly GPU-specific and non-trivial. I recommend http://developer.download.nvidia.com/compute/DevZone/C/html/C/src/reduction/doc/reduction.pdf http://developer.download.nvidia.com/compute/DevZone/C/html/... as an example of how important optimisations lead to those impressive performance numbers.