compute_shaders.rst 14 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380
  1. .. _doc_compute_shaders:
  2. Using compute shaders
  3. =====================
  4. This tutorial will walk you through the process of creating a minimal compute
  5. shader. But first, a bit of background on compute shaders and how they work with
  6. Godot.
  7. .. note::
  8. This tutorial assumes you are familiar with shaders generally. If you are new
  9. to shaders please read :ref:`doc_introduction_to_shaders` and :ref:`your
  10. first shader <toc-your-first-shader>` before proceeding with this tutorial.
  11. A compute shader is a special type of shader program that is orientated towards
  12. general purpose programming. In other words, they are more flexible than vertex
  13. shaders and fragment shaders as they don't have a fixed purpose (i.e.
  14. transforming vertices or writing colors to an image). Unlike fragment shaders
  15. and vertex shaders, compute shaders have very little going on behind the scenes.
  16. The code you write is what the GPU runs and very little else. This can make them
  17. a very useful tool to offload heavy calculations to the GPU.
  18. Now let's get started by creating a short compute shader.
  19. First, in the **external** text editor of your choice, create a new file called
  20. ``compute_example.glsl`` in your project folder. When you write compute shaders
  21. in Godot, you write them in GLSL directly. The Godot shader language is based on
  22. GLSL. If you are familiar with normal shaders in Godot, the syntax below will
  23. look somewhat familiar.
  24. .. note::
  25. Compute shaders can only be used from RenderingDevice-based renderers (the
  26. Forward+ or Mobile renderer). To follow along with this tutorial, ensure that
  27. you are using the Forward+ or Mobile renderer. The setting for which is
  28. located in the top right-hand corner of the editor.
  29. Note that compute shader support is generally poor on mobile devices (due to
  30. driver bugs), even if they are technically supported.
  31. Let's take a look at this compute shader code:
  32. .. code-block:: glsl
  33. #[compute]
  34. #version 450
  35. // Invocations in the (x, y, z) dimension
  36. layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
  37. // A binding to the buffer we create in our script
  38. layout(set = 0, binding = 0, std430) restrict buffer MyDataBuffer {
  39. float data[];
  40. }
  41. my_data_buffer;
  42. // The code we want to execute in each invocation
  43. void main() {
  44. // gl_GlobalInvocationID.x uniquely identifies this invocation across all work groups
  45. my_data_buffer.data[gl_GlobalInvocationID.x] *= 2.0;
  46. }
  47. This code takes an array of floats, multiplies each element by 2 and store the
  48. results back in the buffer array. Now let's look at it line-by-line.
  49. .. code-block:: glsl
  50. #[compute]
  51. #version 450
  52. These two lines communicate two things:
  53. 1. The following code is a compute shader. This is a Godot-specific hint that is needed for the editor to properly import the shader file.
  54. 2. The code is using GLSL version 450.
  55. You should never have to change these two lines for your custom compute shaders.
  56. .. code-block:: glsl
  57. // Invocations in the (x, y, z) dimension
  58. layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;
  59. Next, we communicate the number of invocations to be used in each workgroup.
  60. Invocations are instances of the shader that are running within the same
  61. workgroup. When we launch a compute shader from the CPU, we tell it how many
  62. workgroups to run. Workgroups run in parallel to each other. While running one
  63. workgroup, you cannot access information in another workgroup. However,
  64. invocations in the same workgroup can have some limited access to other invocations.
  65. Think about workgroups and invocations as a giant nested ``for`` loop.
  66. .. code-block:: glsl
  67. for (int x = 0; x < workgroup_size_x; x++) {
  68. for (int y = 0; y < workgroup_size_y; y++) {
  69. for (int z = 0; z < workgroup_size_z; z++) {
  70. // Each workgroup runs independently and in parallel.
  71. for (int local_x = 0; local_x < invocation_size_x; local_x++) {
  72. for (int local_y = 0; local_y < invocation_size_y; local_y++) {
  73. for (int local_z = 0; local_z < invocation_size_z; local_z++) {
  74. // Compute shader runs here.
  75. }
  76. }
  77. }
  78. }
  79. }
  80. }
  81. Workgroups and invocations are an advanced topic. For now, remember that we will
  82. be running two invocations per workgroup.
  83. .. code-block:: glsl
  84. // A binding to the buffer we create in our script
  85. layout(set = 0, binding = 0, std430) restrict buffer MyDataBuffer {
  86. float data[];
  87. }
  88. my_data_buffer;
  89. Here we provide information about the memory that the compute shader will have
  90. access to. The ``layout`` property allows us to tell the shader where to look
  91. for the buffer, we will need to match these ``set`` and ``binding`` positions
  92. from the CPU side later.
  93. The ``restrict`` keyword tells the shader that this buffer is only going to be
  94. accessed from one place in this shader. In other words, we won't bind this
  95. buffer in another ``set`` or ``binding`` index. This is important as it allows
  96. the shader compiler to optimize the shader code. Always use ``restrict`` when
  97. you can.
  98. This is an *unsized* buffer, which means it can be any size. So we need to be
  99. careful not to read from an index larger than the size of the buffer.
  100. .. code-block:: glsl
  101. // The code we want to execute in each invocation
  102. void main() {
  103. // gl_GlobalInvocationID.x uniquely identifies this invocation across all work groups
  104. my_data_buffer.data[gl_GlobalInvocationID.x] *= 2.0;
  105. }
  106. Finally, we write the ``main`` function which is where all the logic happens. We
  107. access a position in the storage buffer using the ``gl_GlobalInvocationID``
  108. built-in variables. ``gl_GlobalInvocationID`` gives you the global unique ID for
  109. the current invocation.
  110. To continue, write the code above into your newly created ``compute_example.glsl``
  111. file.
  112. Create a local RenderingDevice
  113. ------------------------------
  114. To interact with and execute a compute shader, we need a script.
  115. Create a new script in the language of your choice and attach it to any Node
  116. in your scene.
  117. Now to execute our shader we need a local :ref:`class_RenderingDevice`
  118. which can be created using the :ref:`class_RenderingServer`:
  119. .. tabs::
  120. .. code-tab:: gdscript GDScript
  121. # Create a local rendering device.
  122. var rd := RenderingServer.create_local_rendering_device()
  123. .. code-tab:: csharp
  124. // Create a local rendering device.
  125. var rd = RenderingServer.CreateLocalRenderingDevice();
  126. After that, we can load the newly created shader file ``compute_example.glsl``
  127. and create a precompiled version of it using this:
  128. .. tabs::
  129. .. code-tab:: gdscript GDScript
  130. # Load GLSL shader
  131. var shader_file := load("res://compute_example.glsl")
  132. var shader_spirv: RDShaderSPIRV = shader_file.get_spirv()
  133. var shader := rd.shader_create_from_spirv(shader_spirv)
  134. .. code-tab:: csharp
  135. // Load GLSL shader
  136. var shaderFile = GD.Load<RDShaderFile>("res://compute_example.glsl");
  137. var shaderBytecode = shaderFile.GetSpirV();
  138. var shader = rd.ShaderCreateFromSpirV(shaderBytecode);
  139. .. warning::
  140. Local RenderingDevices cannot be debugged using tools such as
  141. `RenderDoc <https://renderdoc.org/>`__.
  142. Provide input data
  143. ------------------
  144. As you might remember, we want to pass an input array to our shader, multiply
  145. each element by 2 and get the results.
  146. We need to create a buffer to pass values to a compute shader. We are dealing
  147. with an array of floats, so we will use a storage buffer for this example. A
  148. storage buffer takes an array of bytes and allows the CPU to transfer data to
  149. and from the GPU.
  150. So let's initialize an array of floats and create a storage buffer:
  151. .. tabs::
  152. .. code-tab:: gdscript GDScript
  153. # Prepare our data. We use floats in the shader, so we need 32 bit.
  154. var input := PackedFloat32Array([1, 2, 3, 4, 5, 6, 7, 8, 9, 10])
  155. var input_bytes := input.to_byte_array()
  156. # Create a storage buffer that can hold our float values.
  157. # Each float has 4 bytes (32 bit) so 10 x 4 = 40 bytes
  158. var buffer := rd.storage_buffer_create(input_bytes.size(), input_bytes)
  159. .. code-tab:: csharp
  160. // Prepare our data. We use floats in the shader, so we need 32 bit.
  161. float[] input = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10];
  162. var inputBytes = new byte[input.Length * sizeof(float)];
  163. Buffer.BlockCopy(input, 0, inputBytes, 0, inputBytes.Length);
  164. // Create a storage buffer that can hold our float values.
  165. // Each float has 4 bytes (32 bit) so 10 x 4 = 40 bytes
  166. var buffer = rd.StorageBufferCreate((uint)inputBytes.Length, inputBytes);
  167. With the buffer in place we need to tell the rendering device to use this
  168. buffer. To do that we will need to create a uniform (like in normal shaders) and
  169. assign it to a uniform set which we can pass to our shader later.
  170. .. tabs::
  171. .. code-tab:: gdscript GDScript
  172. # Create a uniform to assign the buffer to the rendering device
  173. var uniform := RDUniform.new()
  174. uniform.uniform_type = RenderingDevice.UNIFORM_TYPE_STORAGE_BUFFER
  175. uniform.binding = 0 # this needs to match the "binding" in our shader file
  176. uniform.add_id(buffer)
  177. var uniform_set := rd.uniform_set_create([uniform], shader, 0) # the last parameter (the 0) needs to match the "set" in our shader file
  178. .. code-tab:: csharp
  179. // Create a uniform to assign the buffer to the rendering device
  180. var uniform = new RDUniform
  181. {
  182. UniformType = RenderingDevice.UniformType.StorageBuffer,
  183. Binding = 0
  184. };
  185. uniform.AddId(buffer);
  186. var uniformSet = rd.UniformSetCreate([uniform], shader, 0);
  187. Defining a compute pipeline
  188. ---------------------------
  189. The next step is to create a set of instructions our GPU can execute.
  190. We need a pipeline and a compute list for that.
  191. The steps we need to do to compute our result are:
  192. 1. Create a new pipeline.
  193. 2. Begin a list of instructions for our GPU to execute.
  194. 3. Bind our compute list to our pipeline
  195. 4. Bind our buffer uniform to our pipeline
  196. 5. Specify how many workgroups to use
  197. 6. End the list of instructions
  198. .. tabs::
  199. .. code-tab:: gdscript GDScript
  200. # Create a compute pipeline
  201. var pipeline := rd.compute_pipeline_create(shader)
  202. var compute_list := rd.compute_list_begin()
  203. rd.compute_list_bind_compute_pipeline(compute_list, pipeline)
  204. rd.compute_list_bind_uniform_set(compute_list, uniform_set, 0)
  205. rd.compute_list_dispatch(compute_list, 5, 1, 1)
  206. rd.compute_list_end()
  207. .. code-tab:: csharp
  208. // Create a compute pipeline
  209. var pipeline = rd.ComputePipelineCreate(shader);
  210. var computeList = rd.ComputeListBegin();
  211. rd.ComputeListBindComputePipeline(computeList, pipeline);
  212. rd.ComputeListBindUniformSet(computeList, uniformSet, 0);
  213. rd.ComputeListDispatch(computeList, xGroups: 5, yGroups: 1, zGroups: 1);
  214. rd.ComputeListEnd();
  215. Note that we are dispatching the compute shader with 5 work groups in the
  216. X axis, and one in the others. Since we have 2 local invocations in the X axis
  217. (specified in our shader), 10 compute shader invocations will be launched in
  218. total. If you read or write to indices outside of the range of your buffer, you
  219. may access memory outside of your shaders control or parts of other variables
  220. which may cause issues on some hardware.
  221. Execute a compute shader
  222. ------------------------
  223. After all of this we are almost done, but we still need to execute our pipeline.
  224. So far we have only recorded what we would like the GPU to do; we have not
  225. actually run the shader program.
  226. To execute our compute shader we need to submit the pipeline to the GPU and
  227. wait for the execution to finish:
  228. .. tabs::
  229. .. code-tab:: gdscript GDScript
  230. # Submit to GPU and wait for sync
  231. rd.submit()
  232. rd.sync()
  233. .. code-tab:: csharp
  234. // Submit to GPU and wait for sync
  235. rd.Submit();
  236. rd.Sync();
  237. Ideally, you would not call ``sync()`` to synchronize the RenderingDevice right
  238. away as it will cause the CPU to wait for the GPU to finish working. In our
  239. example, we synchronize right away because we want our data available for reading
  240. right away. In general, you will want to wait *at least* 2 or 3 frames before
  241. synchronizing so that the GPU is able to run in parallel with the CPU.
  242. .. warning::
  243. Long computations can cause Windows graphics drivers to "crash" due to
  244. :abbr:`TDR (Timeout Detection and Recovery)` being triggered by Windows.
  245. This is a mechanism that reinitializes the graphics driver after a certain
  246. amount of time has passed without any activity from the graphics driver
  247. (usually 5 to 10 seconds).
  248. Depending on the duration your compute shader takes to execute, you may need
  249. to split it into multiple dispatches to reduce the time each dispatch takes
  250. and reduce the chances of triggering a TDR. Given TDR is time-dependent,
  251. slower GPUs may be more prone to TDRs when running a given compute shader
  252. compared to a faster GPU.
  253. Retrieving results
  254. ------------------
  255. You may have noticed that, in the example shader, we modified the contents of the
  256. storage buffer. In other words, the shader read from our array and stored the data
  257. in the same array again so our results are already there. Let's retrieve
  258. the data and print the results to our console.
  259. .. tabs::
  260. .. code-tab:: gdscript GDScript
  261. # Read back the data from the buffer
  262. var output_bytes := rd.buffer_get_data(buffer)
  263. var output := output_bytes.to_float32_array()
  264. print("Input: ", input)
  265. print("Output: ", output)
  266. .. code-tab:: csharp
  267. // Read back the data from the buffers
  268. var outputBytes = rd.BufferGetData(buffer);
  269. var output = new float[input.Length];
  270. Buffer.BlockCopy(outputBytes, 0, output, 0, outputBytes.Length);
  271. GD.Print("Input: ", string.Join(", ", input));
  272. GD.Print("Output: ", string.Join(", ", output));
  273. With that, you have everything you need to get started working with compute
  274. shaders.
  275. .. seealso::
  276. The demo projects repository contains a
  277. `Compute Shader Heightmap demo <https://github.com/godotengine/godot-demo-projects/tree/master/misc/compute_shader_heightmap>`__
  278. This project performs heightmap image generation on the CPU and
  279. GPU separately, which lets you compare how a similar algorithm can be
  280. implemented in two different ways (with the GPU implementation being faster
  281. in most cases).