NVIDIA Technical Blog

In CUDA PTX, there is an instruction called "shfl.sync.idx", which allows you to shuffle data between threads arbitrarily. However, this shuffling is only efficient if you follow a variety of rules: avoiding bank conflicts and so forth.

Lets say you have 32-threads of a Warp running. Thread0, Thread1, Thread2, etc. etc. Now lets say they all have a variable named "x".

Thread0.x, Thread1.x, Thread2.x, etc. etc.

What "shfl.sync" does, is it allows these "x" variables to be shuffled between all threads in one step.

// Parallel and synchronous: All threads swap
// values arbitrarily. This is a simple "rotate",
// but all combinations are legal for a shuffle

Thread0.x = Thread13.x
Thread1.x = Thread14.x
Thread2.x = Thread15.x
Thread3.x = Thread16.x
...
Thread30.x = Thread11.x
Thread32.x = Thread12.x

That's a swizzle. Now the descriptions on the right? The 15, 16, 17, etc. etc. can be any combination. You can broadcast (have Thread1.x copied to all 32 other threads). You can rotate, you can shuffle. All 32^32 specifications of CUDA-thread transfers is legal.

Now, it may be legal, but broadcasts are slower due to bank conflicts. So following the special rules to make fast swizzles is an annoying task for sure. This paper helps the programmer create optimal swizzles for their program.

For more information:

* https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl-sync

*…

On 1, you are right that most HW most of the time will make the requisite forward progress after the atomic bump. But Vulkan does not guarantee it and most HW doesn't either (you're just getting lucky with a high probability).

The property you're describing is discussed in https://johnwickerson.github.io/papers/forwardprogress_concur2017.pdf, page 1:10:

"After probing our zoo of GPUs with litmus tests, we determined that all GPUs had the following property: once a work-group begins executing a kernel (i.e. the work-group becomes occupant on a hardware resource), it will continue executing until it reaches the end of the kernel. We call this execution model the occupancy bound execution model, because the number of work-groups for which relative forward progress is guaranteed is bound by the hardware resources available for executing work-groups; i.e. the hardware resources determine how many work-groups can be occupant."

However, I know even that property is not true of all GPUs and thus cannot be assumed by Vulkan spec compliant code. (I actually doubt it's true of any GPU's pre-Volta - it's just that the progress needed is often a natural side-effect of switching between threads for latency hiding).

Volta does guarantee this property + more. All starvation-free algorithms are supported by the architecture. See https://devblogs.nvidia.com/inside-volta/. In particular, this line:

"Independent thread scheduling in Volta ensures that even if a thread T0 currently holds…

I don't know where this figure comes from: "Nvidia's official benchmarks show a 10x speed-up when using Caffe with cuDNN.". The graph in the original announcement ( http://devblogs.nvidia.com/parallelforall/accelerate-machine-learning-cudnn-deep-neural-network-library/ ) shows a 11x speed-up for using Caffe on the GPU versus CPU, and a 14x speed-up for using Caffe + cuDNN (again compared to CPU).

Then again these figures are also kind of meaningless, as the specifics of the experiments were not specified. There are a lot of parameters that affect the performance of different implementations in various ways: the number of input feature maps, the number of filters, filter width/height, input width/height. The FFT-approach tends to do well for large filter sizes and lots of input feature maps, for example.

I've played around with cudnn a bit using Theano. The Theano bindings are a work in progress, but so far they've already wrapped the convolution. Compared to conv_gemm (Theano's version of Caffe's GEMM convolution approach), it seems to be sometimes faster, sometimes slower. Soumith Chintala maintains a GitHub repository with benchmarks of various convolution implementations: https://github.com/soumith/convnet-benchmarks
His results aren't that spectacular either.

In my own experiments cuDNN did pretty well for very small filter sizes (i.e. 3x3), often beating even the memory-hungry FFT approach). This is great because the top two scoring entries in the 2014 ImageNet…

The P3 instances are the first widely and easily accessible machines that use the NVIDIA Tesla V100 GPUs. These GPUs are straight up scary in terms of firepower.
To give an understanding of the speed-up compared to the P2 instances for a research project of mine:

+ P2 (K80) with single GPU: ~95 seconds per epoch

+ P3 (V100) with single GPU: ~20 seconds per epoch

Admittedly this isn't exactly fair for either GPU - the K80 cards are straight up ancient now and the Volta isn't sitting at 100% GPU utilization as it burns through the data too quickly ([CUDA kernel, Python] overhead suddenly become major bottlenecks).
This gives you an indication of what a leap this is if you're using GPUs on AWS however.
Oh, and the V100 comes with 16GB of (faster) RAM compared to the K80's 12GB of RAM, so you win there too.

For anyone using the standard set of frameworks (Tensorflow, Keras, PyTorch, Chainer, MXNet, DyNet, DeepLearning4j, ...) this type of speed-up will likely require you to do nothing - except throw more money at the P3 instance :)

If you really want to get into the black magic of speed-ups, these cards also feature full FP16 support, which means you can double your TFLOPS by dropping to FP16 from FP32. You'll run into a million problems during training due to the lower precision but these aren't insurmountable and may well be worth the pain for the additional speed-up / better RAM usage.

- Good overview of Volta's advantages compared to event the recent P100:…

In the paper they say it took 3.14e23 flops to train. They used v100s to do it. This is an insane energy cost (and financial cost).

Nvidia's v100 product page [0] says that it gets about 15 (single precision) - 125 ("deep learning") teraflop/s at 250-300 watts (joules per second). That means that if everything's as perfectly efficient as a marketing product page, it gets about 250/125-300/15 = 2-25 joules per teraflop, putting this model at about 0.6-8 terajoules.

A gallon of gasoline has about 120e6 joules [1] (though if you wanted to compare with burning it in a car, it's only 20-25% efficient at best [2] so it'd be fewer joules/gallon).

This model took the equivalent of about 5,000-67,000 gallons of gasoline at best and at ideal perfect energy efficiency. I get that openAI has made a decision not to be efficient with their dollars in order to see what's possible with future tech, but that means not being efficient with energy either, and it's getting kinda crazy. Sure, microsoft data centers aren't gasoline powered, so maybe it is closer to this ideal energy efficiency, and it's definitely going to be a better carbon footprint, but god damn it just seems wasteful.

Hell, the new A100 (again going off marketing materials [3], so at least it's apples to apples) could do it about 4x more efficiently. Is this research really worth what it costs, when waiting a year makes it that much more efficient?

[0] https://www.nvidia.com/en-us/data-center/v100/

[1]…

Surely, that is a solved problem now that Google has "advance[d] the state of the art in natural language technologies and buil[t] systems that learn to understand and use language in context"? [1]

Or, I guess, one can just buy more NVIDIA GPUs, and simply solve the problem using Torch, following their online guide on "understanding Natural Language with Deep Neural Networks Using Torch" [2].

In general, I hear that many language understanding tasks are finally conclusively conquered, as shown by OpenAI with a combination of supervised and unsupervised learning that has led to an ImageNet-class increase in the state-of-the art results in language understanding datasets [3].

Why, these days, we can even build an "AI System [that] Understands Music Like Humans Do"! (NVIDIA, again; [4])

The inhabitants of Scunthorpe and all the Dicks, Pennys and Dikshits of the world can finally rejoice! Machine learning is on the job!

______________

[1] https://ai.google/research/teams/language/

[2] https://devblogs.nvidia.com/understanding-natural-language-deep-neural-networks-using-torch/

[3] https://blog.openai.com/language-unsupervised/

[4] https://news.developer.nvidia.com/ai-system-understands-music-like-humans-do/