Live data from Hacker News

Nvidia Ampere GA102 GPU Architecture [pdf]

nvidia.com

41–50 of 51 posts

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#41
post #6
post #2

These GPU architectures are too complicated for my taste. It reminds me of the days of segmented memory, and makes me feel like having to jump through hoops all the time. I'm curious if people have tried to invent something more elegant, and where these approaches have failed. Also, the fact that NVidia has probably patented this particular architecture makes it less interesting for me to really dig into.

It's a good question actually. Intel tried to make a GPU called Larrabee that was mostly a bunch of small x86 cores with giant vector units. Turns out that it couldn't compete in rendering performance on existing games (in 2010) without the fixed function units that GPUs have, so they canceled it as a GPU. It did result in the AVX-512 instruction set though. I think the idea still has promise but there's a chicken an…

If I'm reading that right, Doom Eternal only uses compute shader rasterization for writing a lighting acceleration structure where they need to make some fine-grained/coarse-grained decisions depending on depth complexity. The scene is still using rasterization hardware.

Nanite uses compute shader rasterization partly because of the quad overdraw problem since they are targeting near 1 triangle per pixel. But they also say they are using traditional rasterization with recent hardware's addition of mesh shaders when it is faster (which remove a different set of fixed function stuff though, for transform, so still makes the same point).

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#42
post #25

Earlier quoted context omitted.

Its not uncommon for GPU workloads in games to max out about 20% INT32 calculations, but alas its enough to drop the FP32 performance quite a bit. I suspect Nvidia next time will probably separate out the INT32 and 2x FP32 units and gradually move towards going towards a better ratio of hardware that better suits the usual workload split.

Due to the lower amount of INT32 in game loads as you stated, I don't think that separating INT32 and FP32 hardware makes a lot of sense, because you can share a substantial amount of the hardware between the two overall leading to space savings.

On the contrary, "dark silicon" instead suggests that separating fp32 and int32 (now in GA102/104, fp32 and int32/fp32) data paths at the cost of more die space usage currently makes excellent sense. (See also: tensor cores, ray tracing cores.) Jensen Huang very briefly alluded to this when during the GA102/104 announcement he mentioned the end of Dennard scaling.

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#43
post #3

Earlier quoted context omitted.

When you run up against the limits of Moore's law and the end of Dennard Scaling you have to get dirty and do what it takes to get more performance. Personally, I think CPU architecture became too complicated for my taste after the 68k. So what?

You're posting this from a 68k?

Doubt it, but it is possible. I’ve seen a few videos on YouTube of people who’ve gotten old computers (such as the OG Macintosh 128k which runs on 68k) to connect to the internet. I’ve even seen people use slightly “less powerful” machines running on a 6502 (such as the Apple II and Commodore 64).

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#44
post #34

Earlier quoted context omitted.

They did actually make it and it was not supposed to be a GPU..

Sure it was, I was at the GDCE 2009 session on Larrabe.

They showed quake ray tracing demos with it and other people assumed it was supposed to be sold as a GPU instead of listening to what they were actually saying.

No one thought a collection of atom CPUs with AVX512 SIMD was going to be able to compete head to head on rasterization of games with the best Nvidia cards.

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#45

Earlier quoted context omitted.

Fujitsu's A64FX ARM is proof that 512 bit SIMD can work on a CPU based platform.

Sure, but the question is can it render graphics competitively with traditional GPUs.

Probably not, since it's optimized for scientific workloads (being designed specifically for the K computer replacement) (so it doesn't have texture units, ROPs, etc; you'd have to do too much in software to make it actually render things). However I think the overall design is really good and has enormous potential, if not for graphics at the very least for ML.

The vector architectures with extremely high memory bandwidth coming out of Japan recently (NEC SX-Aurora Tsubasa, Fujitsu A64FX) are pretty fascinating.

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#46

Earlier quoted context omitted.

For graphics use GPUs perform a very significant amount of work in hardware (rasterization and texture interpolation being the two computationally most intensive [probably followed by ROP, which blends pixel shader output into the framebufer]; you can easily calculate that the ALU bandwidth of the TMUs is about the same order of magnitude as all the shader cores), which gives them a huge efficiency lead over anything…

Michael Abrash had a great series of articles in Dr. Dobbs detailing how he came to work for Intel (which spun Larabee) after talking at a game conference with some of their people to ask them for a lerp (linear interpolation) instruction in x86 extensions[0] :) Oh and Larabee gave us more than AVX512, it also gave us the Xeon Phis, which were accelerators (much akin to the GPGPU of nvidia GPUs?) aimed at scientific…

OpenCL would have been a bad fit for Xeon Phi.

OpenCL is very specifically tailored for GPUs (though FPGAs may benefit). The concept of "constant memory", "shared memory", and "global memory" is very GPU-centric, and doesn't benefit Xeon Phi at all.

I'd assume that any OpenCL program would simply function better on a GPU, even compared to a 60-core in-order 512-bit SIMD-based processor like Xeon Phi.

---------------

Xeon Phi's main advantage really was running "like any other x86 processor", with 60 cores / 240 threads. But you still needed to AVX512 up your code to really benefit.

Honestly, I think Xeon Phi just needed a few more revisions to figure out itself more. It was on the market for less than 5 years. But I guess it wasn't growing as fast as NVidia or CUDA.

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#47

Earlier quoted context omitted.

Michael Abrash had a great series of articles in Dr. Dobbs detailing how he came to work for Intel (which spun Larabee) after talking at a game conference with some of their people to ask them for a lerp (linear interpolation) instruction in x86 extensions[0] :) Oh and Larabee gave us more than AVX512, it also gave us the Xeon Phis, which were accelerators (much akin to the GPGPU of nvidia GPUs?) aimed at scientific…

OpenCL would have been a bad fit for Xeon Phi. OpenCL is very specifically tailored for GPUs (though FPGAs may benefit). The concept of "constant memory", "shared memory", and "global memory" is very GPU-centric, and doesn't benefit Xeon Phi at all. I'd assume that any OpenCL program would simply function better on a GPU, even compared to a 60-core in-order 512-bit SIMD-based processor like Xeon Phi. ---------------…

Maybe I was mixing up names in my head, but I remember from 5~10 years back an Open[Something] (thought it was OpenCL) that in theory could transparently handle multithreaded code across single/dual/quad[0] core or GPGPU (either nvidia or AMD).

This is what I had in mind when I wrote "if Intel had given it good OpenCL support". Again, maybe I'm mixing things up in my head since my career never took me down that lane to write massively paralell code (though I am a user of it, indirectly, through deep learning frameworks).

[0] back then this was as big a CPU would get

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#48
post #34

Earlier quoted context omitted.

Sure it was, I was at the GDCE 2009 session on Larrabe.

They showed quake ray tracing demos with it and other people assumed it was supposed to be sold as a GPU instead of listening to what they were actually saying. No one thought a collection of atom CPUs with AVX512 SIMD was going to be able to compete head to head on rasterization of games with the best Nvidia cards.

Why would they put texture units on it if it wasn't intended to be sold as a GPU? Consumer gaming GPUs were explicitly planned. Initial released versions even had DirectX drivers. Here is Intel's SIGGRAPH paper featuring benchmarks of Half-Life 2 Episode 2, Gears of War, and F.E.A.R. http://download-software.intel.com/sites/default/files/m/9/4...

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#49

Earlier quoted context omitted.

OpenCL would have been a bad fit for Xeon Phi. OpenCL is very specifically tailored for GPUs (though FPGAs may benefit). The concept of "constant memory", "shared memory", and "global memory" is very GPU-centric, and doesn't benefit Xeon Phi at all. I'd assume that any OpenCL program would simply function better on a GPU, even compared to a 60-core in-order 512-bit SIMD-based processor like Xeon Phi. ---------------…

Maybe I was mixing up names in my head, but I remember from 5~10 years back an Open[Something] (thought it was OpenCL) that in theory could transparently handle multithreaded code across single/dual/quad[0] core or GPGPU (either nvidia or AMD). This is what I had in mind when I wrote "if Intel had given it good OpenCL support". Again, maybe I'm mixing things up in my head since my career never took me down that lane…

There's a version of OpenCL that compiled to Intel, but I'm not very familiar with it.

I remember reading things like: https://software.intel.com/content/www/us/en/develop/documen...

Where you'd have to use float8 types to be assured of SIMD-benefits on CPU code. As such, its probably more useful to rely upon auto-vectorizers in C++ code (such as #pragma omp simd) and maybe intrinsics for the complicated cases.

Intel does seem to have some level of OpenCL -> AVX tech: http://llvm.org/devmtg/2011-11/Rotem_IntelOpenCLSDKVectorize...

Re: Nvidia Ampere GA102 GPU Architecture [pdf]

#50
post #42
post #25

Earlier quoted context omitted.

Due to the lower amount of INT32 in game loads as you stated, I don't think that separating INT32 and FP32 hardware makes a lot of sense, because you can share a substantial amount of the hardware between the two overall leading to space savings.

On the contrary, "dark silicon" instead suggests that separating fp32 and int32 (now in GA102/104, fp32 and int32/fp32) data paths at the cost of more die space usage currently makes excellent sense. (See also: tensor cores, ray tracing cores.) Jensen Huang very briefly alluded to this when during the GA102/104 announcement he mentioned the end of Dennard scaling.

But the GA102/GA104 doesn’t have seperate execution units for INT and FP32 because the INT also does FP32. So I don’t see how that shows that separating FP32 and INT hardware makes sense.
Post reply on HN