NVIDIA Technical Blog

I think you just pick a problem you want to solve with gpu programming and go for it. Learning what you need along the way. Nvidia blog posts are great for learning things along the way such as https://devblogs.nvidia.com/cuda-pro-tip-write-flexible-kernels-grid-s…

You could just use LoopVectorization on the CPU side. It's been shown to match well-tuned C++ BLAS implementations, for example with the pure Julia Gaius.jl (https://github.com/MasonProtter/Gaius.jl), so you can follow that as an example for getting BLAS-speed CPU side kernels. For the GPU side, there's CUDAnative.jl and KernelAbstractions.jl, and indeed benchmarks from NVIDIA show that it at least rivals directly writing CUDA (https://devblogs.nvidia.com/gpu-computing-julia-programming-language/)… so you won't be missing anything just by learning Julia and sticking to using just Julia for researching new kernel implementations.

If you're interested in some more details: https://devblogs.nvidia.com/vulkan-raytracing/ Do note that the following is part of the article and it's from 2018, not sure about it's accuracy anymore: "As denoted by the ‘NVX’ prefix, this API is not yet final and could undergo some minor changes before the final release."

Also worth checking out OmniSci (formerly MapD), which features an LLVM query compiler to gain large speedups executing SQL on both CPU and GPU: https://github.com/omnisci/omniscidb . And here's a link to a blog post giving a high level overview of the advantages of JIT compilation of queries over an interpreter: https://devblogs.nvidia.com/mapd-massive-throughput-database-queries-l… .

Disclosure: I work on Google Cloud (even with Zak sometimes). The DGX-2h is a beast! Don’t parse this as “huh, TPU Pods are about the same as just a few V100s”. The data sheet [1] is probably the easiest to follow, but their writeup is more informative [2]. These are souped up V100s, with awesome networking, which is pretty similar in style to a TPU Pod. So I’d say that they’re both purpose built systems for distributed ML training. The name for the NVIDIA system is even “DGX SuperPOD” :). [1] https://www.nvidia.com/content/dam/en-zz/es_em/Solutions/Data-Center/d… [2] https://devblogs.nvidia.com/dgx-superpod-world-record-supercomputing-e…

In short, they created a tool that makes it easier to optimize for GPUs. Currently you need a good grip of- https://devblogs.nvidia.com/register-cache-warp-cuda/ to optimize your pipeline which is not pedestrian. Now you can define what you want to do, and let Swizzle Inventor come up with the best way to exploit the register cache, apparently as well as top CUDA engineers. Great for real-time computer vision applications that might need to solve large nonlinear least square problems.

>That's not true. There's a lot of heavily serial algorithms which do not always work well on the GPU unless they are of a very large size, larger than most applications. This happens a lot in optimization and differential equations, two of the pillars of scientific computing. And those situations tend to get a very minimal speedup from SIMD. >Tim already wrote that article where it shows how Julia code generation can be faster in some instances for GPUs. https://devblogs.nvidia.com/gpu-computing-julia-programming-... Yes, that's great, and would have made for an interesting challenge, rather than focusing on SIMD instead of serial or GPU performance. >Where? I didn't word it well, but I meant to say that Julia doesn't have much overhead compared to C/C++, unlike Python (what overhead it does have comes from garbage collection).

Someone should make a discrete gpu with an arm cpu core or 50 sharing the same memory and maybe ssd. Then you have a really nice hierarchy of parallelism, 2 fast arm cores, 10 medium, 50 slow, then gpu. Would be great for stuff like multigrid methods where you are balancing parallelism and latency: https://devblogs.nvidia.com/parallelforall/high-performance-geometric-…

We have much, much more coming soon, in addition to LIDAR/radar data. You can join in with the fun here: https://udacity.com/self-driving-car It's worth pointing out some other awesome datasets, such as: • http://research.comma.ai • https://devblogs.nvidia.com/parallelforall/deep-learning-self-driving-… • http://data.selfracingcars.com Any questions I can answer?

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]…

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…

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… *…

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… [3] https://blog.openai.com/language-unsupervised/ [4] https://news.developer.nvidia.com/ai-system-understands-music-like-hum…

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:…

The article misses the core issue: neural network architectures are still in flux. "Neuromorphic" chips are hardwired to one architecture, which makes them power efficient but less flexible. When designs are more stable, such chips could be more practical. Meanwhile, upcoming (non-neuromorphic) AI processors are taking two directions: larger numbers of simplified GPU-type cores (such as NVIDIA Xavier and Intel's Lake Crest/Nervana chips), and FPGAs. Simplifying cores means lower precision, as fp32 and fp64 are overkill for neural networks and take up lots of silicon. The current NVIDIA Pascal added fp16 and byte operations such as the DP4A convolution instruction[1]. Even smaller precision is practical (down to 1 bit with XNORnet[2], and the DoReFa paper[3] gives an excellent summary of the falloff in accuracy through 32-8-4-2-1 bits for weights, activations, and gradients). [1] https://devblogs.nvidia.com/parallelforall/mixed-precision-programming… [2] XNORnet, https://arxiv.org/abs/1603.05279 [3] DoReFa, https://arxiv.org/abs/1606.06160

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-… ) 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…

Have you looked into using DALI [1] to do augmentation on GPU? They've gotten some nice speedups for computer vision that way. [1] https://devblogs.nvidia.com/fast-ai-data-preprocessing-with-nvidia-dal…

No, the A100 has a 19.5 TFLOP theoretical peak for SGEMM[1], real world benchmarks will likely achieve 93% of that, and so the DGX A100 will be 145 TFlops of FP32 SGEMM performance or 0.145 FP32 PFLOPS. Maybe in 72 FP64 TFLOPS. FP64 is what the TOP500 benchmaks count.[2] The 5 "petaflops" number is a creatively constructed marketing number based on FP16 TensorCore "flops", sparse matrix calculations, and then multiplying by 8x for some reason. They basically take the 19.5 FP32 TFLOPS number and multiply it by 32x to get to the claimed 624 "TFLOPS" for a single A100. 8 * 624 = 5 "petaflops". I see they get 2x by actually using FP16 instead of FP32, 2x from counting sparse matrix ops as dense ops, and 8x from somewhere else that I have no idea. [1] https://devblogs.nvidia.com/nvidia-ampere-architecture-in-depth/ [2] https://www.top500.org/resources/frequently-asked-questions/

Even more details: https://devblogs.nvidia.com/nvidia-ampere-architecture-in-depth/

The official docs are pretty good. This is probably a good starting point: https://devblogs.nvidia.com/even-easier-introduction-cuda/ Nvidia also sponsored a Udacity course, but I don't have any direct experience with it: https://developer.nvidia.com/udacity-cs344-intro-parallel-programming

Easily outperformed by the more traditional TensorRT, which TF also supports: https://devblogs.nvidia.com/tensorrt-integration-speeds-tensorflow-inf… In fact, also seems to be outperformed by plain PyTorch using a single V100: https://github.com/NVIDIA/DeepLearningExamples/tree/master/PyTorch/Cla…

Looks like the GPU stores collections of triangles then used to perform the ray computations. https://devblogs.nvidia.com/practical-real-time-ray-tracing-rtx/

If you dont know already: Azure has started investing heavily in offloading the task from CPU to GPU and FPGA. There are really good framework on .NET to achieve the same. NVIDIA CUDA GPU - https://devblogs.nvidia.com/hybridizer-csharp/ GPU scripting and CUDA - http://www.aleagpu.com/release/3_0_4/doc/ FPGA - https://hastlayer.com/project, https://netfpga.org/site/#/

I was trying to understand the significance of this. There is a blog post here[0]. (I had a brief glance, please correct me if I'm wrong) It seems to be a framework allowing neural networks (and layers of neural networks) to be composed together more easily. It's not completely clear to me what the advantages are without digging in, I assume to simplify/standardise the interface between the layers? [0] https://devblogs.nvidia.com/neural-modules-for-speech-language-models/

Even if GPT-2 were released, very very few would have the hardware to run it because of gpu ram running out (and doing some sort of load-unload system would make training times unfeasibly long). And those who have the hardware to run it, has probably already made a version of their own or reasons not to. So I'm wondering if this GPT-2 hype is a genuine concern of openai, or if it's mostly a PR flex to say 'Look at us, we made a good model!'. As an example, look here by Nvidia https://devblogs.nvidia.com/training-bert-with-gpus/ who made GPT-2 8B, which is ~5 times as large as GPT-2.

You are literally as wrong as it is possible to be wrong. http://hoxxep.github.io/webgl-ray-tracing-demo/ https://devblogs.nvidia.com/accelerated-ray-tracing-cuda/

This NVIDIA post goes into extending Bayesian Optimization to multiple metrics [0]. It shows how you can use efficient optimization to find a good Pareto Frontier[1]. [0]: https://devblogs.nvidia.com/sigopt-deep-learning-hyperparameter-optimi… [1]: https://en.wikipedia.org/wiki/Pareto_efficiency

Seems to lose in Mobilenet and win in other benches according to Nvidia's benchmarks. I'll wait for independent testing before I drop $100. https://devblogs.nvidia.com/jetson-nano-ai-computing/

https://developer.nvidia.com/embedded/buy/jetson-nano-devkit https://devblogs.nvidia.com/jetson-nano-ai-computing/

That's really cool! The only other optimization that I wanted to try and do was to use GPU acceleration to support huge numbers of shapes and entities. Something similar to what's described here [0]. I think it would really help for very very large numbers of entities or large complex geometry. [0] - https://devblogs.nvidia.com/thinking-parallel-part-i-collision-detecti…