16-bit Floating-point Error and Activation Function Tolerance

In the last post we showed that small error in the activation function used for a network did not negatively impact the performance or output of a neural network (using our relatively small networks).

As part of our work on systems which are error tolerant, we took the different activation functions that are commonly used and an approximation known as the “serpentine” curve and ran it through our networks to see how they performed to determine if error in the activation function or the shape of the function had much of an impact.

Below we show the series of results for each function with a 16-bit implementation and after that there is a table showing the amount of error relative to a full-precision for each.

The error is calculated by taking every half-precision floating point value between 0.0 and 5 and passing it into the target function and measuring the different from the full 64-bit implementation, as opposed to the earlier work where we compared approximations to their source.

From this we can see that dropping to 16-bit floating point doesn’t incur any penalty to the functions with only linear component (Relu, LinearX, Constant) as these functions all take an input that exists in half-precision and return the same input or a fixed output that is also a valid half-precision number without any chance of overflow or rounding.

On the other hand our ‘tanh’ and ‘serpentine’ functions both perform calculations in half-float leading to errors which propagate to the return value. We can see by the median value that most results do no incur very much error at all but some values are very wrong.

ErrorChart.png

Noticeably the approximation of the ‘Tanh’ function, ‘Serpentine’, has lower maximum and average error. To inspect this further we plotted the error for every input in the range 0.0 to 5.0 for both functions.

In both results the inherent pattern of the floating-point numbers is visible, but much noticeable on the ‘serpentine’ plot. The exact pattern is a little worrying as it changes the pattern from being close enough to ‘random’ to possibly having some implication in error propagation.

In our tests the network we used appears to be tolerant enough of either error and give good enough results for both - however the sigmoid-like ‘serpentine’ function is much cheaper than the ‘tanh’ which raises the question why it isn’t being used instead.

Additionally, the error across the range is non-linear for both. If we were to use this in a network the weights which allow for input closer to 0 will have a more diverse set of choices for values and allowing for a higher-precision fit than those with much larger activation values. This should result in a change in how the network converges on the correct answer.

More importantly is the propagation of error. In neural networks there is the problem of exploding and diminishing gradients which is commonly known (and one of the reasons for the popularity of non-linear but basically linear functions like ReLu). Allowing error to propogate through an application means that the combined error could cause unwanted behaviour such as that with gradients but also that the gradient you are using during back propagation may not be correct for the activation function that is being used.

This leaves us with the question of how tolerant are networks to small errors in the activation function, particularly with low precision data-types. Are there certain patterns of input or weights which in conjunction with non-linear error would slow or distort the learning process? Or will this type of error on larger networks limit the convergence onto the best possible solution?

Sometimes Activation Functions Only Vaguely Matter

This is a follow on from the last post where we discussed improving neural network performance by allowing for cheaper approximate functions to be used instead of the costly full-fat activation functions.

In Machine Learning an activation function is used to decide which neurons should be activated. It is a transform of the signal from layer to the next. To be able to fit to complex functions an activation layer must be non-linear otherwise the network is only capable of simple linear regression.

A list of common activation functions can be found on the wikipedia page which also gives the properties and limitations of each.

As you can see in that table, there is a lot and they seem to be very strictly defined. If we were to stick with ‘tanh’ as we looked at in the last post we could see it being implemented with more or less error. We were interested in figuring out if an implementation of ‘tanh’ which allowed an enormous amount of error, but was still non-linear, would be unable to perform the same task as the standard implementation of the ‘tanh’ function.

So we came up with two alternatives. ‘ApproxTanh3’ an order 3 polynomial approximation and ‘ApproxTanh12’ an order 12 polynomial approximation. The plot of these against the standard implementation can be seen below:

approximations.png

As you can see, the high order approximation is a reasonable fit to the actual tanh implementation where as the low order approximation basically falls off a cliff around x=0.5 .

So, what happens when we use this in a real neural network?

For this task we have chosen the standard MNIST handwriting dataset and a pretty simple convolutional neural network.

The dataset is a very common one, it is a collection of 28x28 pixel images of handwritten numbers and the number they represent. The standard test error rate of a convolutional neural network on this dataset is between ~0.2% and ~1.7% for different standard implementations.

And our neural network looks like this where {CHOSEN ACTIVATION} is either one of our approximations or the real ‘tanh’ implementation:

Conv2D(32, (5, 5), input_shape=(1, 28, 28), activation='relu')
MaxPooling2D(pool_size=(2, 2))
Dropout(0.2)
Flatten()
Dense(128, activation = {CHOSEN ACTIVATION} )
Dense(num_classes, activation='softmax')

With this network with the same seed we get these results:

Function TEST ERROR RATE (%)
Keras.Tanh 0.97
ApproxTanh3 0.84
ApproxTanh12 1.31

This was very surprising to us. We expected the terrible order 3 approximation to break the ability for the neural network to fit the problem, but it’s results were within the ranges we expected and beat the standard implementations error rate in this specific seed (Other runs show some variance of +/- one percent).

In this example, our terrible activation function was able to match the standard implementation. It seems for certain problems where the inputs to the activation stay within certain ranges that the exact shape of the activation function isn’t important as long as it is vaguely correct within the working range.

This means that there is a large range of error tolerance in this system. A fact which it seems, for some networks, could be exploited to reduce learning times!

Improving the Performance of TensorFlow Activation Functions

In our research we are currently investigating existing applications which are tolerant to functions which return a value that is within a tolerance range rather than relying on absolute unit-in-last-place guarantees.

Machine-learning is a popular area of computing at the moment which just happens to conform to this paradigm. From specialised GPUs to TPUs we can see large teams approaching machine learning with variable precision stages and noise inducing layers to speed up or aid in the learning process.

While there is a lot of areas here which show levels of error tolerance, we have chosen to focus on the activation functions. In particular we are looking at the popular ‘tanh’ activation.

tanh33.png

The function ‘tanh’ is a useful function due to its non-linearity and convenient output mapping of [-1,1]. Unfortunately, Tanh is an expensive function to compute accurately.

So, if we want to improve this we need to begin with the facts we know already:

  • Some activation layers in a neural networks can be expensive

  • Machine learning algorithms often use low-precision floats (16/32-bit)

  • Floating-point cannot give absolutely correct answers for real numbers

  • Inputs to some learning algorithms are often normalised to be values between zero and one

  • Accuracy between zero and one in floating-point is non-linear. So error is higher in functions working closer to one

  • Approximations of complex functions can result in better performance

  • Machine learning is inherently error-tolerant

  • Machine learning has a major problem with run-times being too long

From these facts we can assume:

  1. The current ‘tanh’ function implementation has an acceptable and variable level of error

  2. An approximation of ‘tanh’ which has better performance and acceptable levels of error will not impact on the overall outcome of the learning process and be valuable to the programmer

  3. Any implementation which does not impact the overall result negatively and improves run-time would be a positive contribution.

So, if we want to show our assumptions apply to the real-world, we need to prove each one.

1) This is the simplest assumption to prove. The Tensorflow website provides a few machine learning tutorials, one of which is on Text Classification. This example uses a sigmoid activation in it’s final step. Changing this to a ‘tanh’ activation has no negative impact on the learning process. The accuracy of the output model remains in the 80-90% accuracy range.

Since we are using ‘tanh’ and the output is acceptable, the error that we know it has due to its underlying type must be valid and hasn’t impacted the convergence of the learning process.

2) Next we want to take an approximation and compare the result of using that to the official implementation to see if it has any difference characteristics which might be detrimental.

To do this we must select an implementation for the replacement. We have kept to a simple polynomial for this example as the resulting error is very small on the limited input range we are using ([0,1]). Although there are other, better performing, approximations with different costs associated.

With our replacement Activation function selected, we simple replace the final activation layer to point to our function and can then run the learning process. To ensure fairness between all implementations we use a fixed seed and reset the global state of Tensorflow for each run. This will allow us to see the ways in which the implementations may diverge through the learning process.

ML_TanHActiviations.png


The tests give us the above results. The first chart is the standard implementation, while the other three are order-16 polynomials implementations in 64, 32 and 16-byte floating-point. As can be seen, there is very little difference between the approximations. The approximations do diverge from the official implementation around epochs 15-20. This divergence is expected due to the accumulation of small differences between the implementation but has no impact on the overall accuracy of the resulting model. In all tests that were ran the approximations and official implementation all converged to around the same accuracy.

This shows that the approximations do not negatively impact the overall result of the model.

NOTE: These tests were all trained on a CPU due to the NVIDIA libraries being non-deterministic. This non-determinism prevented 1-to-1 comparison of different activation functions with the same seed.

3) Next we need to consider the performance impact. As we are limited to working with a very small model due to problems with NVIDIA’s implementation preventing 1-to-1 testing on the GPU the activation layer is only a small part of a small model and this makes it difficult to accurately measure through the noise. To get around this we will be testing the implementations on the CPU separately (until we do further work with more time!).

Run on (8 X 3600 MHz CPU s)
CPU Caches: L1 Data 32K     (x4),  L1 Instruction 32K   (x4),
            L2 Unified 262K (x4),  L3 Unified 8388K     (x1)
------------------------------------------------------
Benchmark                                      Time  
------------------------------------------------------
Tanh Standard Implementation            |      212 ms 
                                        |
--Order 16 Polynomial Implementations-- |
SSE           64bit Float               |      114 ms 
Nonvectorised 64bit Float               |      144 ms 
SSE           32bit Float               |      114 ms 
Nonvectorised 32bit Float               |      129 ms 

Different machine learning libraries will take different approaches to vectorisation. As our functions are trivially vectorised on the CPU we see huge improvements over the standard ‘tanh’ performance, but we also see large improvements even in the standard scalar implementations.

A win of 25-50% on performance is significant when running a network with hours spent training, even if only a portion of that is spent with the activation functions.

Conclusion

With these three assumptions shown to be true in limited practice we are in a good position to make stronger assertions for larger neural networks and hopefully present a method to improve the performance of your Tensorflow models without having to make any tangible sacrifice!

Installing Tensor Flow for GPU

The TensorFlow website provides a brief outline of how to get a tensorflow programs running on an NVidia card but due to updates on the NVidia site, some of the instructions are a little hard to follow and it is easy to make a mistake, so I will give an updated guide here.

STEPS

  • Run DDU, a tool for removing current display drivers from your computer.
    DDU removes the current NVidia driver from your computer and any associated files. This prevents the a common error when rolling back to older NVidia drivers. 
  • Install the CUDA Toolkit version 9.0. This toolkit comes with the correct driver version for using it. There are newer versions of the toolkit now, but version 9.0 is what is required for TensorFlow-GPU. If the link here breaks, this version of the toolkit is found in the archives.
  • Install cuDNN v7.0.5 (Dec 5, 2017), for CUDA 9.0. Like the CUDA Toolkit there are newer versions of this tool, so this one is found in the archives. 

The TensorFlow website suggests to set the correct PATH variables in Windows according to the documentation - but in my experience with these specific installers that is done during the installation process.

Once everything is installed and you have restarted your computer you can install tensorflow-gpu from pip or through Anaconda

pip3 install --upgrade tensorflow-gpu

Then you can run the test python scripts found here to validate the GPU support. 
NOTE: There is a rather nasty bug in the example GPU scripts - when it suggests to test a maximum percentage usage of memory on the GPU ensure that the maximum number you set is less than what is currently available. Any contention for memory on the GPU causes a hard driver crash.

With that all up and running - enjoy much faster machine learning!

What can we learn from GPU Frame Captures: Stardew Valley

If you read my last post on frame captures we discussed how GPU frame captures can be used to analyse performance, and determine how a scene is rendered for a specific game. 

It also gives insight into what the game is doing right (from a performance perspective) and how it can be improved.

In the last post we looked into Europa Universalis 4 where we found an astronomical amount of draw-calls, strange scaling behaviour and sub-optimal texture packing to render a relatively simple (although quite large) world.

In this post we are looking at a retro game for comparison. Many people who havent kept up with modern graphics development may believe that modern "retro" looking games use the same old rendering techniques that used to be mandatory on the old consoles - this isn't the case!

Stardew valley mimics the aesthetic of old SNES games like Harvest Moon but the process by which it draws the scene makes it more dynamic, smooth and pretty doing many things that would have been impossible back then to make a great experience now.

Rendering Pattern

The rendering pattern is very simple for this 2D game. It follows the expected patternn of drawing the floor and background followed by a distance sorted order of objects. In a 2D game with this view, the objects "further away" are those higher-up the screen as they will not obscure those further down. This allows the renderer to not have to worry about depth testing.

The full pattern for rendering as I could figure it out from a capture appears to be:

track.png
stardewrendering.png

In this pipeline the only steps which really take any time are the two full screen passes, they are the initial ground rendering after the screen clear, and the full-screen weather effects- in our case, snow.

The initial ground rendering is a simple tiled geometry representing the game world grid and each grid cell has UV coordinates to map to a cell in the bound ground textures. 

The snow rendering is similar but mapping to a different section of the texture as shown below:

The snow is then animated by updating which cell is read for each game-world quad each frame. Notice that the background for the snow texture is not black, it is a dark grey. This is what adds the hazy effect to the scene as the texture is not directly written to the scene as the other blocks have been, it is an additive effect (D3DBLENDOP_ADD) resulting in softening of the image.

All in all it produces a quite nice, and retro effect - even though this approach might be too intensive for most real "retro" games.

snowComparison.png

Another interesting feature of this game is that it lets you customise how your character looks. Instead of using those options to generate a new sprite sheet for rendering your character, the game instead keeps all the possible options in memory and builds up the character chunk by chunk based on your choice from these options, each frame.

clothing.png

This is a neat and simple trick that avoids having to create a new texture when a game is loaded - but limits the extension and number of choices for the player as adding any new options increases the memory usage of the game through the entire run-time, although this is probably trivial on a game of this scale.

So how do we improve on this?

A game of this scale and already tiny frame rendering times (<2ms) there isnt really any need to improve anything. Unless you were running an incredibly outdated machine this game will run incredibly well.

However, if we were thinking of running this game on a very limited device we would probably want to consider the techniques that were used on the actual old retro games.

For starters, the majority of the scene (with an exception of the weather effects) doesn't change frame to frame - and those bits which do are somewhat predictable. There is no reason why parts of the scene couldnt be rewritten as things or the camera move and otherwise kept the same - this would infact throw out a majority of the draw-calls when the camera is stationary. The draw backs to this approach is that any mistakes in the implementation have really ugly artefacts.

This game also falls prey to the same issues as EU4 when it comes to texture sizes (although most are trivial in this case making the slight error not too big of a problem).  There are a number of large textures in this game that are just slightly above a power of two texture size - resulting in some wasted memory.

Otherwise, this frame capture is incredibly nice and easy to investigate and navigate.

What can we learn from GPU Frame Captures: Europa Universalis 4

A common technique used in the games industry to analyse the state of a game in development and look for where improvements can be made is to capture a frame(all the GPU work being done between each frame being shown to the user) from the game and analyse the results to see the impact of the different elements on the screen and how they might be able to be rearranged.

In doing this we can look at the relative costs of the different aspects of the processes in the scene. An example of this would be that the programmer suspects that the newly implemented bloom post-effect has had too big an impact on performance. So the programmer grabs a few frames from the game and looks at the time that is being taken to perform that effect at different places in the game, and what the cost of that is relative to the rest of the scene.

Another use for this approach is for new programmers coming onto a complex project to get a quick look at the rendering 'pipeline' that has been implemented. As in, what is rendered when, by what shader and in what order. This is quite useful for someone who only needs to make a minor change on a project.

An interesting side-effect of this is that we are able to frame-capture fully finished games and look at what the code is asking our computer to do and from that derive how the rendering system of that game works to some extent.

In this post we will be looking at Europa Universalis 4(EU4) frame captures and constructing a flow chart of how that game is rendered. EU4 is the forth game in the Europa series from Paradox Interactive released in 2013.

To capture frames from the game we will be using Intel's GPA. GPA is one of the simpler and less detailed frame capturing tools available but is good enough for this example and will allow someone with GPU to be able to follow along if they wish.

(For those interested in more complex capture tools the most popular amongst those I know in industry is RenderDoc, but AMD and NVidia each have there own tools which provide specific functionality for features of their cards in Radeon GPU Profiler and NVIDIA NSight respectively. Microsoft also have a tool called PIX but I haven't personally used that in some time but have heard it works with AMDs tools now.)

GraphicsMonitorStartPage.png

First thing we want to do is use the GPA Graphics Monitor tool to launch EU4 from its directory. This will launch the game with the GPA overlay, giving us real-time performance information and the option to capture frames. In this example I will be running on a NVIDIA 1080 GPU so will be targeting the max settings so that we can see where each option is placed in the pipeline. The settings and overlay can be seen in the screenshot below.

GraphicsMonitorSettings2.png

The intention was to run with max settings, but we have had to step down from 4k to 2560x1440 due to the game performance dropping to below 1FPS at the higher resolution while running with Intel GPA (and not good without it either...) . Something that will be covered later.

With the settings covered and GPA running we now need to take some frame captures. As this game has a relatively simple view, we will be taking a capture close to the ground, mid-zoom out and full zoom-out to get a full coverage of the some of the most common cases.

This is performed by pressing Ctrl+Shift+C when the game is in the right position. An important note when capturing is to take note of the CPU usage. This data is not always captured with the frame and can indicate if there is a CPU bottleneck which may be causing poor GPU performance. In the case of our capture (below) we arent using very much CPU at all so that should not impact the results much here.

In my captures I have used a save file from a playthrough with a random 'new world'. This is due to my own interest in seeing if the randomly generated terrain is handled differently than the standard world map geometry.

Once the three captures have been taken, open Intels Graphics Frame Analyser tool where the captures should be displayed.

frameanalyser.png

We can start by opening the zoomed-out capture first, as this has the least detail and should give use the broad strokes of the rendering process. This will give you a view of the data like this:

faroutcapture.png

The important things to look at here are the top timeline and the view frame at the bottom left. The timeline on the top will allow you to set the x and y axis to be the GPU duration to highlight particularly expensive calls and the view on the bottom will highlight what is being rendered on each render target for each of the calls in the timeline.

simpleselection.png

As an example I selected the group of similar draw calls (ID=~1000 to ~10400) which is appears to handle only the drawing of the borders in of the world and this is shown in the bottom frame with a pink highlight of the pixels which are being rendered to. Above that there is information on my selection. It says that we selected 9483 draw calls which account for 90.1%(!!!) of the total rendering time for the frame! Scary, but we will get to that in the section on how to improve this pipeline. For now this is just to show you how this tool works.

We will use this approach to see how the scene is rendered by stepping through the draw calls and seeing what they are doing. From what we can gather the process for drawing the world appears to be starting with the zoomed out capture:

Zoomed-out capture analysis
 

 With the numbers referring to the GPU instruction ID

With the numbers referring to the GPU instruction ID

timelinefar.png

The timeline is broken up by some taller commands, there are the buffer copies which nicely split the pipeline into the four coloured sections shown in the diagram. The length of each section is not entirely representative of the total rendering time but is an indicator that something isn't as ideal as it could be in this particular pipeline.

Overall the pipeline is relatively simple (by modern standards). It uses a depth map for correctly layering objects in the scene and each stage of pipeline ties nicely into the options from the settings. With the exception of shadows which appear to be disabled or not rendering at this level of zoom. 

From a graphics programmers view, there is a number of very interestingly strange things going on with the way objects are submitted for drawing in this pipeline.

For starters the sheer number of draw calls is ridiculously large for even a major AAA game and this leads to a bottleneck in the gpu as each of the submitted draw calls are relatively tiny with only tens of primitives at the best of times. This is shown, rather crudely, in GPA by the occupancy and thread dispatch boxes showing red for the frame.

occupancyproblem.png

This is starving the GPU. For a lot of these calls the GPU is probably spending more time in the driver than it is actually processing the data it is being given. This is particularly noticed when the game is rendering borders, state names, rivers and trees.

The terrain is also interesting. The world is rendered as 36, 8196-triangle grids. Essentially one big height map broken into chunks, which reads its lighting from the precomputed base texture in the setup stage. 8196-triangles is a relatively insignificant amount of triangles to draw and each of the 36 grids appears to share the same texture and shader state. There doesn't appear to be an apparent reason this is separated instead of being one giant geometry. If the reason is LOD (level of detail) related this could be resolved CPU side very cheaply to select and combine multiple vertex buffers or vertex buffer sections into one call.

There is a lot of other details but they appear to be the same throughout all captures, so I will address them in the optimisation section at the end.

Close Capture Analysis

shadowmap.png

As suspected,. the shadows showed up in the close zoom levels. In the zoomed out pipeline map there was some strange depth maps being cleared and appeared not to be used later in the pipeline. When we have a closer view these render targets become used for storing and processing the shadow map that is produced in the setup stages, the shaders being used later in the pipeline are also binding this as a texture and appear to be using it - the rest of the pipeline remains unchanged in my casual exploration.

EU4_RenderingClose.png

The borders, text and UI remain as the main culprit of GPU use, but with the additional of a lot of calls rendering the scene objects for the shadow map and processing adding some extra work.

This is offset by some good culling of the terrain. In the far out LOD the entire terrain and water were rendered as many chunks together. At this level, the fewer chunks reduces the total number of draw calls submitted, but the culling was not correct and some out of view chunks were submitted.

timelinenear.png

Mid-Zoom Capture

Our mid-level capture appears to have the same pattern as the far-out zoom. All the city, unit detail and shadow is LOD'ed out, leaving only the map rendering. The gives us the same behaviour with less of the draw-calls as less is visible. Rendering borders still covers the majority of the calls, but it isnt as dramatic as the far out zoom.

midlevel.png

We can see in this image that the depth buffers used for the shadow map are still bound and cleared each frame (RT2-5) and are apparently unused, but I may have missed something.

Sins of a Render Empire

So in the last three sections I gave an overview of the pipeline and how it changes in different configurations - essentially some features are disabled as LOD features based on distance. However, this isnt the only place to look when we are considering performance. So in this section we are going to cover a few oddities that need addressing if this game was going to be optimised.

4k Support

As I mentioned in the setup for this, I intended to run this on max setting at 4k resolution. This currently isn't possible with a high-end intel chip and NVidia 1080 - a little strange for a game from 2013.

The game runs roughly fine at 4k when we dont enable all the on-off options in the video menu. A quick investigation into this and it seemed that the shadows really take most the time. A problematic quirk of resolutions is that when we double the size we quadruple the number pixels (which also quadruples the amount of work). In this engine, the shadow buffers appear to be sized based on the full rendering resolution. So, going from 1440p to 4k doesnt just quadruple the cost of rendering the scene, it has to be rendered twice at that size, so it is a ~8x increase in rendering cost. 

Additionally the trees are expensive at 4k. This is a relatively simple reason, at 4k we see more pixels on each tree. Each tree has high texture detail and for some reason a detailed normal map. So now we have also quadrupled the texture read cost, with little coherency because the trees are small and dense and the texture resolution is reasonably high.

Water is similarly effected, but not as much as the trees due to similarity in the pixel space helping cache coherency.

In the game there are a number of what appear to be generated textures, the size of these appears to be loosely based on the resolution that is being used. However, they do not stick to power of two texture sizes. So a texture that is 1029x1029  is actually a 2048x2048 texture under the hood on certain hardware, this doesnt effect the appearance of the texture but does have performance implications and is just a massive waste of memory.

This next complaint is just because I play this game a lot. At 4k the menus get really tiny and that's a pain.

DirectX 9.0c

This game is written with DirectX 9.0c. In 2013 when this came out, DirectX 10 and 11 had long been a standard and DirectX 12 was well on the way as well as AMD's early experiments with Mantle which led to Vulkan.

AMD and NVIDIA put a lot of work into optimising drivers for modern hardware. This focus is obviously for patterns and use-cases in common software. DirectX 9.0c misses a lot of features that could really make this type of game fly on even a basic laptop.

Overbinding and draw calls

Textures are bound each time they are used. Tiny objects are submitted to draw calls. There is either a massive lack of batching or it is not apparent in the capture.

This puts the game at the mercy of the PCI-bus and the drivers. Every time you ask the GPU to do something you risk a major state change stalling the next draw. So much of the textures being bound are identical (outside of the UI) and there is a lot of texture slots to bind, some textures are so low resolution they may as well be constant buffers and allow for some nice out of order operations to be able to be done.

Texture Space Wastage

The cubemap for the sea is a good example of this. Due to the angle of the camera only a limited section of the cubemap will be accessed, so half of the cube map is just blank. This is more likely a trade-off than an error as the DX9 cubemaps have some strange rules - but it isn't ideal in a modern game.

cubemap.png

 

Conclusion

From this we have shown how to analyse a game through frame capture. Shown how to extract the pattern it uses for rendering and view the content of each rendering step. We then covered what could be causing some of the performance hiccups we see during gameplay and how they could be fixed.

I would like to add the captures frames to this post but I think that might count as distributing content from the game and I am not sure of the legality of that, so for now I hope my instructions on how to set this up on your own system with your own copy of the game is enough!

 

C++17: The new problem with 'auto'

Since C++11 introduced 'auto' there have been discussions about whether it increases or decreases the readability and comprehensibility of code. 

Personally, I believe that auto is useful for making code concise when used with an IDE that can resolve the type for you when needed without having to go through too much digging around, but can be harmful is overly used or applied in non-obvious ways.

C++14 extended on the use of 'auto' in a logical fashion. If it is alright to use for type definitions then it should be acceptable as the return type in the definition of a function where type can be deduced from the return statement or a trailing definition added to the end of the function definition.

This I find to be a little bit less reasonable in the first case as it demands the programmer to actively explore the implementation of a function to understand its usage, but due to the limitations there can only be one return type so finding and understanding any single path through the function will give you a thorough understanding of the type that is being returned.

This isn't too terrible, but it leaves the user experience of the programmer being a little unnecessarily tedious as the first half of the function definition is very unclear.

Lets look to the other half of the function declaration, the inputs to a function. From C we already had variadic inputs which allow for any number of inputs to a function, this was then further extended to to variadic function templates in C++11. 

In combination with the function auto return type we now have a function that can take any arguments and return any single-type that. This exact types that are in play or acceptable are non-obvious from the function definition and require full understanding of the source implementation to be able to use safely. This is very bad and the current state of play as of C++14 - but not as bad as it is going to get.

Now, to the point of this post. C++17 introduces the very much sought after compile-time if statement in the form of 'if constexpr(...)'. This allows for whole blocks of function to be discarded or included based on a logical check at compile-time. Very very useful and a great addition that could simplify a lot of code and produces more efficient output by giving more information to the compiler.

However, if we consider alongside what we have been discussing so far we will see that this changes the behaviour of the function auto return type. Where as in earlier versions of C++ the auto would refer to a single return type (unless some complicated templating was in use) we can now have a function of arbitrary return type based on a compile-time decision. Changing our single deduced return type with arbitrary input into an arbitrary return type with arbitrary inputs. Essentially removing all useful information from the definition of the function and requiring a full understanding of all control paths through the function to fully know which inputs are valid and what it will return.

This is a problem of weakly typed languages and one of the strengths of C++ was not having this problem. It leads to very confusing code like this:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
//Abusive Case:
template<class... Args>
auto AutoFunction(Args...args)
{
    constexpr int n = sizeof...(Args);
    auto argtuple = std::make_tuple(args...);

    if constexpr (n == 1)
    {
        if constexpr(std::is_same<type_list<Args...>::type<0>, int>::value)
        {
            return 0;
        }

        else if constexpr(std::is_same<type_list<Args...>::type<0>, float>::value)
        {
            return 0.f;
        }

    }
    else if constexpr (n == 2)
    {
        if constexpr    (   std::is_same<type_list<Args...>::type<0>, float>::value
                        &&  std::is_same<type_list<Args...>::type<1>, float>::value
                        )
        {
            return false;
        }
    }
    else
    {
        return std::map<void*, int>();
    }

}

int main()
{
    //All cases.
    auto a = AutoFunction(2.f);        //returns float
    auto b = AutoFunction(2);          //returns int
    auto c = AutoFunction(2.f, 2.f);   //returns bool
    auto d = AutoFunction();           //returns std::map

    auto input = SomeFunction(12);
    auto whatAmI = AutoFunction(input);

    return 0;
}

In this example 'AutoFunction' is essentially acting as four different functions and which function it is behaving as will be determined by the result-type of 'SomeFunction' which itself could have the same problem.

The number of lines of code needed to be able to correctly and safely use 'whatAmI' has went from simply the definition of AutoFunction to the entire function as well as any functions which may feed as the input.

This is a terrible way to be able to acceptably write code. From the outside the function appears to be sensible but can hide strange behaviour. Programmers are far from constantly vigilant and this will only lead to problems.

What is especially problematic with this way of writing code is that it is actually very powerful. There are numerous algorithms and patterns which could be improved this way and may result in a better compiled output. It is simply that the behaviour is not clear, it is not signaled that it may behave that way and therein lies our problem.

I don't want 'if constexpr' removed, it is incredibly useful. I don't want 'auto' return types removed either. I simply believe that for them to be a non-dangerous addition there needs to be something else present to make the programmer using the function aware.

 

 

 

 

 

Automated Function Replacement (Short Quick Post)

Ok! Quick short post.

There are many ways to write even trivial functions. For example, f(x) = (5.f * x) can be written in C++ in ways that will produce different output assembly. Shown here is the Compiler Explorer view of a few of those functions:

Notice that despite them all being simple only a few implementations are transformed to the single 'mul' operation. This is even when compiled with the '-fp:fast' fast-math flag which removes the restrictions which should allow this transformation.

Since small operations like this are common in most applications this is a problem effecting most programs that are being built with these modern C++ compilers. (Disclaimer: Intel seemed to handle this better).

Our latest research into function generation includes a step that can resolve this problem. when passed a function such "Five_b" above, it can correctly identify at build time using specific machine learning approaches that this is simply "f(x) = (5.f * x)" and return the source code to replace it (admittedly, pretty ugly at the moment).

More interestingly we can identify sub-expressions of larger functions and return replacements or approximations of those sub-expressions.

A very impressive result of the system so far is that it is able to recognise complex functions such as the sum of integers below 'x' and replace with the simple form of "f(x) = (x(x+1))/2". The same can be done for replacing Fibonacci sequences with approximations, and replacing very expensive functions with table look-ups.

The current training set is limited, but each test adds to our database and improves the efficiency of the system.

We will be using this system going forward for automated approximate function generation and to aid in our exploration of the accuracy-performance design space.

CUDA Minimal Setup

As in the OpenCL post, the default samples that are shipped with the CUDA SDK are a big mess of complicated. (Although some of the online resources are better). As such here is a minimal implementation of the same simple setup of the most basic things in GPGPU. 

CUDA is a little different than OpenCL. In C++ if you aren't separately compiling and linking, it is written like it is part of the C++ language and those parts of the code are compiled with NVCC.

If you want to build a CUDA application in Visual Studio the easiest way is to create a new project from the Visual Studio home screen and select NVIDIA CUDA, alternatively you can switch the compilation for each cpp file in your solution browser to use the NVIDIA compiler. This should all be available if you installed the CUDA Toolkit (SDK) 

cudamenu.png

With all that being said, here is the simple demo doing the same as the OpenCL demo. That is: Initialise the device, allocate some memory, run a kernel to fill that memory with 42, finish running the kernel, copy the data back to the host CPU and check it is valid.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
#include "cuda_runtime.h"

__global__ void DoSomething(int *_inData)
{
	//This gets us the threadid.
	const unsigned int threadIndex = threadIdx.x;
	_inData[threadIndex] = 42;
}

#define SAMPLE_COUNT 128
void StartCuda()
{
	//Allocate host and device memory
	int size = sizeof(int) * SAMPLE_COUNT;
	int* hostBufferMemory = new int[SAMPLE_COUNT];
	int* cudaBufferMemory;
	cudaMalloc((void **)&cudaBufferMemory, size);

	//Run the kernel
	int num_threads = SAMPLE_COUNT;
	dim3 grid(1, 1, 1);
	dim3 threads(num_threads, 1, 1);
	DoSomething<<<grid,threads>>>(cudaBufferMemory);// << < grid, threads >> > ();
	if (cudaSuccess != cudaGetLastError())	return;

	//Wait for work to finish
	cudaDeviceSynchronize();

	//Copy Buffer to host (CPU)
	cudaMemcpy(hostBufferMemory, cudaBufferMemory, size, cudaMemcpyDeviceToHost);
	if (cudaSuccess != cudaGetLastError())	return;

	//Check our magic number was set.
	if (hostBufferMemory[6] != 42)
		return;

	//Release memory because we are being well behaved.
	cudaFree(cudaBufferMemory);
	delete[] hostBufferMemory;

	return;
}

int main()
{
	//Find CUDA Devices and set the first valid one we find.
	int deviceCount;
	cudaGetDeviceCount(&deviceCount);

	if (deviceCount == 0)	return 0;
	else					cudaSetDevice(deviceCount - 1);

	StartCuda();
}

OpenCL Minimum Setup

If you look at the sample code available for using OpenCL 1.2/2.0 on from the primary vendors you will notice that it is all very complicated, some is out of date ( i.e using depreciated functions) or difficult to get running.

The demos they provide do some clever things and definitely show you some good tricks but they are far from a good entry point. Ideally a newcomer to GPGPU wants to be able to open a single file, compilable demo. It should be simple and cover the basics only. Nothing more.

The user should be able to walk from the top of the file to the bottom without having to jump around and see the working pipeline in-order and clearly. Once they are comfortable with seeing the minimum and can have a play around, then you can show them more than just the top layer.

If those of us in the high-performance computing end of software development havent learned anything (and sometimes I think we havent learned anything) from javascript, GUI development tools and the rapid pace of the app development world we should hopefully have at least learned that getting someone onto your platform and working quickly is the best way to keep them. A single hour to understanding how to use something basically is better than a whole day of stress to only gain slightly more.

But enough ranting. I have written a minimal code demo for OpenCL in this style. It lacks all the options, robustness, safety and control the of Intel samples - but its basically ~100 lines of code instead of many thousand and is enough to show the basic concepts and usage patterns.

  1
  2
  3
  4
  5
  6
  7
  8
  9
 10
 11
 12
 13
 14
 15
 16
 17
 18
 19
 20
 21
 22
 23
 24
 25
 26
 27
 28
 29
 30
 31
 32
 33
 34
 35
 36
 37
 38
 39
 40
 41
 42
 43
 44
 45
 46
 47
 48
 49
 50
 51
 52
 53
 54
 55
 56
 57
 58
 59
 60
 61
 62
 63
 64
 65
 66
 67
 68
 69
 70
 71
 72
 73
 74
 75
 76
 77
 78
 79
 80
 81
 82
 83
 84
 85
 86
 87
 88
 89
 90
 91
 92
 93
 94
 95
 96
 97
 98
 99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
//Include the OpenCL Headers
// On Intel default directory is: C:\Intel\OpenCL\sdk\include\CL
// During installation of the Intel SDK it usually placed $(INTELOCLSDKROOT) into your
// system environment. So the project should have the default include directory of: $(INTELOCLSDKROOT)\include
// On AMD and NVIDIA this is different. When you grab the SDK from their site it will usually tell me
// but it is mostly variations on this. Easiest way to find it if the site it confusing is to open a sample
// in visual studio and check the include directories listed in the project properties.
#include <CL/cl.h>

// Standard library to make some things easier.
#include <vector>
#include <string>
#include <fstream>
#include <streambuf>

#define BUFFER_ENTRY_COUNT 256

int main()
{
    //The important objects for initialising and using OpenCL.
    cl_platform_id		platform	= 0;
    cl_device_id		device		= 0;
    cl_context			context		= 0;
    cl_command_queue	queue		= 0;

    //First get our platform-------------------
    {
        //Find which platforms are available
        cl_uint numPlatforms = 0;
        cl_int err = clGetPlatformIDs(0, 0, &numPlatforms);
        if (err != CL_SUCCESS || numPlatforms == 0) return 0;

        //Fetch the IDs and take our first one.
        std::vector<cl_platform_id> availablePlatforms(numPlatforms);
        err = clGetPlatformIDs(numPlatforms, &availablePlatforms[0], 0);
        if (err != CL_SUCCESS) return 0;
        platform = availablePlatforms[0];
    }

    //Now we need our device------------------
    {
        //You can specify if you want CPU/GPU or Accelerator here, but for simple getting going
        //we will just take any.
        cl_device_type  deviceType = CL_DEVICE_TYPE_ALL;

        // Same as above, get the number of devices before we fetch the information.
        cl_uint num_of_devices	= 0;
        cl_int err				= clGetDeviceIDs(platform, deviceType, 0, 0, &num_of_devices	);
        if (err != CL_SUCCESS || num_of_devices==0) return 0;

        //Fetch the ids and select the first one.
        std::vector<cl_device_id> deviceVector(num_of_devices);
        err	= clGetDeviceIDs(platform, deviceType, num_of_devices,	&deviceVector[0], 0);
        if (err != CL_SUCCESS) return 0;
        device = deviceVector[0];
    }

    //Create the context minimal code.
    {
        cl_context_properties contextProps[3];
        contextProps[0] = CL_CONTEXT_PLATFORM;
        contextProps[1] = cl_context_properties(platform);
        contextProps[2] = 0;

        cl_int err = 0;
        context = clCreateContext(&contextProps[0], 1, &device, 0, 0, &err);
        if (err != CL_SUCCESS) return 0;
    }

    //Create a Queue
    {
        cl_int err = 0;
        cl_command_queue_properties props= 0;
        queue = clCreateCommandQueueWithProperties(context, device, &props, &err);
        if (err != CL_SUCCESS) return 0;
    }

    std::string CLProgramFilename	= "./simpleprogram.cl";
    std::string CLProgramKernelName = "EmptyKernel";
    std::string CLProgramSource = "";
    cl_program  CLProgram = 0;
    cl_kernel   CLProgramKernel = 0;
    //Read program source code from file
    {
        std::ifstream file(CLProgramFilename);
        std::string temp;
        while (std::getline(file, temp)) 
        {
            CLProgramSource.append(temp);
        }
    }

    //Create Program from source
    {
        //Take the source and get the program
        cl_int err;
        const char* text = CLProgramSource.c_str();
        cl_program program = clCreateProgramWithSource(context, 1, &text, 0, &err);
        if (err != CL_SUCCESS) return 0;

        //Build it for your specified device.
        err = clBuildProgram(program, (cl_uint)1, &device, "", 0, 0);
        if (err != CL_SUCCESS) return 0;

        //Pull out the kernel(function) we want to use from the program.
        //Programs can have many kernels
        CLProgramKernel = clCreateKernel(program, CLProgramKernelName.c_str(), &err);
        if (err != CL_SUCCESS) return 0;
    }

    cl_mem outputBuffer = 0;
    cl_uint buffSize = BUFFER_ENTRY_COUNT * sizeof(cl_int);
    //Create an output Buffer
    {
        //We are creating a buffer here. The important flags are the CL_MEM_... ones.
        // In this example we say we want one that the kernel can only write and the CPU
        // can only request to read.
        // There are many options here and combining them in different ways has interesting performance effects.
        cl_int	err		= 0;
        outputBuffer	= clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_HOST_READ_ONLY, buffSize, NULL, &err);
        if (err != CL_SUCCESS) return 0;
    }

    //Run Kernel
    {
        //Set the buffer we write to. This maps to the index of the variable in the function in the kernel.
        cl_int err = clSetKernelArg(CLProgramKernel, 0, sizeof(cl_mem), (void *)&outputBuffer);

        //Global size is the total number of things we want to do.
        //Local size is the chunks we are breaking it into. If global not divisible by local
        //it will throw an error.
        cl_uint globalSize	= BUFFER_ENTRY_COUNT;
        cl_uint localSize	= 16;
        err = clEnqueueNDRangeKernel(queue, CLProgramKernel, 1, NULL, &globalSize, &localSize, 0, NULL, NULL);
        if (err != CL_SUCCESS) return 0;

        //Ensuring all the work is done before we copy out our buffer to check the kernel ran correctly.
        err = clFinish(queue);
        if (err != CL_SUCCESS) return 0;

        //Validate the output from our buffer
        cl_int ourOutput[BUFFER_ENTRY_COUNT];
        err = clEnqueueReadBuffer(queue, outputBuffer, CL_TRUE, 0, buffSize, ourOutput, 0, NULL, NULL);
        if (err != CL_SUCCESS) return 0;

        //Check the array has the magic number in it
        if (ourOutput[6] != 42)	return 0;
    }

    //Everything went well.
    return 0;
}

I hope this is useful to anyone trying to get started with OpenCL. I will do the same for Vulkan Compute and DirectX11/12 if I see this gets any traction.