Как стать GPU-инженером за час

Современный мир не был бы таким, какой он есть, без GPU-вычислений. Современные консольные игры, VR, AR, криптовалюты, машинное обучение — всё это работает на горячих графических процессорах.

Однако среди мобильных разработчиков видеокарты не пользуются большой популярностью: многие думают, что это очень сложно, а некоторые вообще не замечают, что iPhone в принципе имеет видеокарту.

Этим докладом хотелось бы ознакомить широкие массы разработчиков с программированием графических процессоров, с прицелом на мобильные платформы и, конечно же, трендовые темы.

Transcript

1. Master GPU-engineering in one hour Andrey Volodin Senior Matrix Multiplicator
2. Agenda • Computer graphics history • Modern rendering • Apple side of things • What is GPGPU? • Metal Compute Shaders • Hype train 2
3. History 3
4. 1977 Empire Strikes Back (Atari 2600) The first notable system to use sprite graphics 4
5. 1977 Atari 2600 128 bytes of RAM including call stack and the state of game world • Typical resolution: 160x192 • 128 colors in palette • 160 * 192 * 7 bits = 26 880 bytes per frame No framebuffer, graphics were generated in real-time. Literally. 5
12. 1977 Atari 2600 VCS could only display five interactive objects at any one time: • 2 «player» sprites • 2 «missile» sprites • 1 «ball» sprite 12
13. «Racing the beam» the VCS could only display five interactive objects at any one time: two "player" sprites, two "missile" sprites, and one «ball», but once the electron beam had drawn a sprite, the program could shift the position of said sprite horizontally and redraw it 13
14. «Racing the beam» the VCS could only display five interactive objects at any one time: two "player" sprites, two "missile" sprites, and one «ball», but once the electron beam had drawn a sprite, the program could shift the position of said sprite horizontally and redraw it 14
15. «Racing the beam» the VCS could only display five interactive objects at any one time: two "player" sprites, two "missile" sprites, and one «ball», but once the electron beam had drawn a sprite, the program could shift the position of said sprite horizontally and redraw it 15
16. Blind spots are the only times programmers could do anything that didn't involve drawing graphics on the screen, such as computing joystick inputs, player movements, scoring 16
17. 1977 Nintendo Entertainment System 1983 (Vice Project Doom) 17
18. 1977 Nintendo Entertainment System 1983 • 8-bit colors • Still no framebuffer • PPU (Picture Processing Unit) • Tiled graphics (aka «character graphics») • Operates with tiles of 8x8 (or 8x16) pixels • 8 sprites per scanline • Another advantage is collision detection 
 (movable/non-movable sprites) 18
19. 1977 Nintendo Entertainment System 1983 (Vice Project Doom) 19
21. 1977 Second Generation - Shaded Solids 1983 1987 1991 … • Very expensive, mostly used in professional simulators • Vertex lighting • Rasterization of filled polygons • Depth buffer and blending 21
22. 1977 Second Generation - Shaded Solids 1983 1987 1991 … 22
23. 1977 First «GPU» 1983 1987 1991 … 1999 NVidia releases first «Graphics Processing Unit» - GeForce 256 • Made a definition of what GPU should be • Achieved 10 million polygons processed in 1 second • Vertex transform • Lighting • Barely programmable 23
24. 1977 GeForce3 with GeForceFX first programmable GPU 1983 1987 1991 … 1999 2001 • Introduced a concept of shaders • Vertex and fragment operations • Macro assembly language • Very limited ADDR R0.xyz, eyePosition.xyzx, -f[TEX0].xyzx; DP3R R0.w, R0.xyzx, R0.xyzx; RSQR R0.w, R0.w; MULR R0.xyz, R0.w, R0.xyzx; ADDR R1.xyz, lightPosition.xyzx, -f[TEX0].xyzx; DP3R R0.w, R1.xyzx, R1.xyzx; RSQR R0.w, R0.w; MADR R0.xyz, R0.w, R1.xyzx, R0.xyzx; MULR R1.xyz, R0.w, R1.xyzx; DP3R R0.w, R1.xyzx, f[TEX1].xyzx; MAXR R0.w, R0.w, {0}.x; 24
25. Recent trends 25
26. Recent trends Time Trans MHz GFLOPS Aug02 121M 500 8 Jan03 130M 475 20 Dec03 222M 400 53 • 1.8x increase of transistors • 20% decrease in clock speed • 6.6x GFLOP speedup 26
27. Modern rendering process 27
28. • GPUs are very limited in what they can do • Can only draw primitives: triangles, lines, points • Highly optimized for floating point operations 28
29. 3D objects are stored as a set of triangles 29
30. They are placed in the 3D scene which is simply a coordinate space 30
31. Next, there is usually a camera to provide different view angles 31
32. Last step in vertex stage: projection 32
33. Affine transform matrices Same concepts in UIKit (CGAffineTransform) Translation Scale Rotation (Y) 33
34. Affine transform matrices x x x 34
35. Typical vertex shader return uniforms.modelViewProjectionTransform * float4(newPosition, 1.0); 35
36. Typical vertex shader return uniforms.modelViewProjectionTransform * float4(newPosition, 1.0); Calculated on CPU as a combination of all transforms 36
37. Next, rasterizer comes into play 37
38. Next, rasterizer comes into play 38
39. Fragment shaders Two triangles, that cover the screen 39
40. Fragment shaders fragColor = vec4(1.0, 0.0, 0.0, 1.0); 40
41. Fragment shaders fragColor = vec4(length(deltaToCenter), 0.0, 0.0, 1.0); 41
42. Fragment shaders sin(time) 1.0 -1.0 42
43. Fragment shaders fragColor = vec4(length(deltaToCenter) * sin(time), 0.0, 0.0, 1.0); 43
44. Fragment shaders 66 LOC 140 LOC 800 LOC 44
45. Fragment shaders Business card challenge by Paul Heckbert (1984) 45
46. Custom shader example Imagine we have a sphere 46
47. Custom shader example What if we will displace every vertex by a smooth noise value? 47
50. Custom shader example Next, we will take gradient texture and will read as high as current displacement is 50
51. Custom shader example And also offset read coordinates every frame 51
52. Custom shader example 52
53. Draw calls happen for every set of geometry with unique render state 53
54. 54 Draw calls happen for every set of geometry with unique render state
55. 55 Draw calls happen for every set of geometry with unique render state
56. What Apple has to offer? 56
57. 2007 OpenGL ES 1.1 (iPhone 2G) 57
59. 2007 OpenGL ES 1.1 (iPhone 2G) 2010 OpenGL ES 2.0 (iPhone 4) 59
61. 2007 OpenGL ES 1.1 (iPhone 2G) 2010 OpenGL ES 2.0 (iPhone 4) 2016 OpenGL ES 3.0 (iPhone 7) 61
62. GLKit.framework OpenGLES.framework SpriteKit SceneKit Cocos2D-ObjC libGDX Unity Unreal Cocos2D-X 62
63. Announced by Khronos in 2015 Initially, Apple was part of the working group 63
65. Low CPU overhead Modern GPU features Do expensive tasks less often Optimized for CPU behaviour Thinnest possible API 65
67. Low CPU overhead Modern GPU features Do expensive tasks less often Optimized for CPU behaviour Thinest possible API 67
68. Low CPU overhead Modern GPU features Do expensive tasks less often Optimized for CPU behaviour Thinest possible API 68
69. Low CPU overhead Modern GPU features Do expensive tasks less often Optimized for CPU behaviour Thinest possible API 69
71. Low CPU overhead Modern GPU features Do expensive tasks less often Optimized for CPU behaviour Thinest possible API 71
72. MTLDevice 72
73. MTLDevice Represents a single GPU 73
74. MTLDevice MTLCreateSystemDefaultDevice() NOTE: Don’t treat this object as a singletone 74
75. MTLDevice MTLCommandQueue 75
76. MTLDevice MTLCommandQueue Retrieved from device via: device.makeCommandQueue() 76
77. MTLDevice MTLCommandQueue Often referred to as «Metal context» 77
78. MTLDevice MTLCommandQueue 78
79. MTLDevice MTLCommandQueue = MTLCommandBuffer 79
80. MTLDevice MTLCommandQueue = MTLCommandBuffer guard let commandBuffer = commandQueue.makeCommandBuffer() else { fatalError("Could not create command buffer") } Made on per-queue basis: 80
81. Command types: 1. Render commands 2. Blit commands 3. Compute commands 81
82. 1. Render commands 2. Blit commands 3. Compute commands MTLRenderCommandEncoder MTLBlitCommandEncoder MTLComputeCommandEncoder 82
84. MTLRenderCommandEncoder commandBuffer.makeRenderCommandEncoder(:) 84
85. MTLRenderCommandEncoder commandBuffer.makeRenderCommandEncoder(:) 85
86. Pipeline state • Represents GPU state that is need to be set for the current command • Must be initialized with shader functions • Each pipeline state type has its own optional parameters • Usually being cached and reused 86
87. Pipeline state // Create a reusable pipeline state for rendering geometry let stateDescriptor = MTLRenderPipelineDescriptor() stateDescriptor.vertexFunction = vertexFunc stateDescriptor.fragmentFunction = fragmentFunc 87
88. Pipeline state let ps = try device.makeRenderPipelineState(descriptor: stateDescriptor) 88
89. MTLRenderCommandEncoder MTLRenderPipelineState .setRenderPipelineState(pipeLineState) 89
90. MTLRenderCommandEncoder MTLRenderPipelineState .setVertexBuffer(:index:) Geometry buffer 90
91. MTLRenderCommandEncoder MTLRenderPipelineState .setVertexBuffer(:index:) Geometry buffer Uniform buffer .setFragmentBuffer(:index:) Texture 0 91
92. MTLRenderCommandEncoder .drawPrimitives(***) +1 92
93. MTLRenderCommandEncoder .drawPrimitives(***) +1 Another state Another geometry +1 93
94. MTLRenderCommandEncoder .endEncoding() +1 +1 94
95. +1 +1 // Send buffer to the command queue commandBuffer.commit() // Wait until all command are executed commandBuffer.waitUntilCompleted() // Subscribe to completion event commandBuffer.addCompletionHandler {} 95
96. Metal Compute Shaders 96
97. Early GPGPU 1999-2001 • Hoff (1999): Voronoi diagrams on NVIDIA TNT2 • Larsen &McAllister (2001): first GPU matrix multiplication (8-bit) • Rumpf & Strzodka (2001): first GPU PDEs (diffusion, image segmentation) • NVIDIA SDK Game of Life, Shallow Water (Greg James, 2001) 97
98. Early GPGPU 1999-2001 • PHD in computer graphics to do this • Financial companies hired game developers 98
99. 2002 GPGPU.org General Purposed GPU 1999-2001 99
100. 2002 1999-2001 100
101. 2002 1999-2001 R G B A R G B A 101
102. 2002 1999-2001 R G B A R G B A 0.17 0.21 0.1 0.2 0.1 0.21 0.2 0.0 102
103. 2002 1999-2001 2007 CUDA • First GPU arch. and software platform designed for computing • First C/C++ language and compiler for GPUs • 2007 began a massive surge in GPGPU development 103
104. 2002 1999-2001 2007 CUDA Output registers Output registers Thread ID Input registers Fragment program Thread program 104
105. Metal Compute Shaders • Act just like fragment or vertex shader, but general purposed • Programmed with keyword kernel • Suitable for highly parallel tasks • Can be put in the same command buffer with render/blit commands 105
106. Task: multiply every element in float buffer by a certain value Purely parallel thing - suitable for compute shaders 106
107. Code Time! 107
108. 1. Declare class ArrayProcessor 2. Use MTLDevice or MTLCommandQueue as a dependency injection 3. Cache static elements in init(:) 108
109. public class ArrayProcessor { public let commandQueue: MTLCommandQueue public let device: MTLDevice public let bufferMultiplierPipelineState: MTLComputePipelineState public init(commandQueue: MTLCommandQueue) { … } … } 109
110. Next, implement encoding GPU work on CPU side 4. Prepare type-container for kernel’s parameters fileprivate struct Uniforms { public let multiplier: Float public let count: UInt32 } NOTE: Be careful with Swift’s memory layout, use C/C++ delcarations to avoid tricky bugs 110
111. 5. Encode compute kernel command into command queue public class ArrayProcessor { … public func process(array: [Float], multiplier: Float) { … } … } MTLDevice MTLCommandQueue MTLComputePipelineState 111
112. public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } MTLDevice MTLCommandQueue MTLComputePipelineState 112
113. public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } MTLDevice MTLCommandQueue MTLComputePipelineState Array buffer Uniform buffer 113
114. public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } MTLDevice MTLCommandQueue MTLComputePipelineState Array buffer Uniform buffer 114
115. public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } MTLDevice MTLCommandQueue MTLComputePipelineState Array buffer Uniform buffer MTLComputeCommandEncoder 115
116. MTLDevice MTLCommandQueue MTLComputePipelineState Array buffer Uniform buffer MTLComputeCommandEncoder public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } 116
117. MTLDevice MTLCommandQueue MTLComputePipelineState Array buffer Uniform buffer MTLComputeCommandEncoder public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } 117
118. Threads and threadgroups • Metal executes your kernel function over 1D, 2D or 3D grid • Each point in the grid represents a single instance of your kernel function • That is called thread • Threads are organized together into threadgroups that can share common block of memory 118
119. Threads and threadgroups 119
120. Threads and threadgroups kernel void myKernel(uint2 threadgroup_position_in_grid [[ threadgroup_position_in_grid ]], uint2 thread_position_in_threadgroup [[ thread_position_in_threadgroup ]], uint2 threads_per_threadgroup [[ threads_per_threadgroup ]]) 120
121. Threads and threadgroups kernel void myKernel(uint2 threadgroup_position_in_grid [[ threadgroup_position_in_grid ]], uint2 thread_position_in_threadgroup [[ thread_position_in_threadgroup ]], uint2 threads_per_threadgroup [[ threads_per_threadgroup ]]) 121
122. Threads and threadgroups 122
123. Threads and threadgroups 123
124. Threads and threadgroups Threads in a threadgroup are executed in SIMD way (Single Instruction Multiple Data) if All threads execute both branches, keep divergence to minimum 124
125. Threads and threadgroups The division of threadgroups into SIMD groups is defined by Metal SIMD group size is returned by threadExecutionWidth of compute pipeline state object All you have to do is define threadgroup size 125
126. 6. Calculate threadgroup count and size et executionWidth = bufferMultiplierPipelineState.threadExecutionWidth et threadgroupsPerGrid = MTLSize(width: (buffer.count + executionWidth - 1) / executionWidth, height: 1, depth: 1) et threadsPerThreadgroup = MTLSize(width: executionWidth, height: 1, depth: 1) 126
127. MTLDevice MTLCommandQueue MTLComputePipelineState Array buffer Uniform buffer MTLComputeCommandEncoder public class ArrayProcessor { public func process(array: [Float], multiplier: Float) { } } computeEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadsPerThreadgroup) 127
128. 7. Write shaders Write your kernels kernel void bufferMultiplier(device float* inputBuffer [[buffer(0)]], const device BufferMultiplierUniforms& uniforms [[buffer(1)]], const uint threadIndex [[ thread_position_in_grid ]]) { if (threadIndex >= uniforms.bufferSize) { return; } const float initialValue = inputBuffer[threadIndex]; inputBuffer[threadIndex] = initialValue * uniforms.multiplier; } 128
129. Benchmarks What we will be playing with: var inputBuffer = [Float](repeating: 1.0, count: 1_000_000) let multiplier: Float = 2.0 What we will be comparing to: // CPU Implementation for i in 0..<inputBuffer.count { inputBuffer[i] = inputBuffer[i] * multiplier } 129
130. • Metal finished in 0.006s • CPU finished in 0.1s • Which is ~17 times slower Benchmarks 130
131. 1_000_000 1_000 Benchmarks 131
132. • Metal finished in 0.003s • CPU finished in 0.0001s • Which is ~30 times faster Benchmarks 132
133. Tips 1. Beware of memory alignment 2. Beware of CPU-side encoding overhead 3. Keep code divergence to minimum 4. Use half instead of float whenever possible 5. Avoid using ints 6. Calculate threadgroup sizes thoughtfully 7. Cache reusable CPU-side objects 8. Don’t wait for GPU to finish execution 133
134. Metal Performance Shaders 134 • A framework of data-parallel algorithms for the GPU • Optimized for iOS • As simple as calling a library function
135. Metal Performance Shaders 135 Lanczos Resampling
 Histogram, Equalization, and Specification Median Filter Thresholding Image Integral iOS9
136. Metal Performance Shaders 136 iOS10 MPSCNN
137. Metal Performance Shaders 137
138. Metal NN Graph API 138
139. CoreML 139
140. CoreML 140
141. CoreML 141 • Easy to use • Wide range of desktop frameworks • Almost as fast as manual encoding • GPU/CPU optimizations • Is not customizable • Sometimes buggy • Zero control
142. CoreML 142 • Easy to use • Wide range of desktop frameworks • Almost as fast as manual encoding • GPU/CPU optimizations • Is not customizable • Sometimes buggy • Zero control DEPRECATED
143. CoreML 143 Freshly baked: 1. 16-bit floats 2. Tensorflow-Lite converter 3. Custom layers (both CPU and GPU)
144. CoreML 144 • Easy to use • Wide range of desktop frameworks • Almost as fast as manual encoding • GPU/CPU optimizations • Is customizable • Still a bit buggy • Zero control
145. Thanks! @s1ddok @s1ddok siddok@gmail.com 145
146. @s1ddok 146 https://medium.com/@s1ddok