Attention: Here be dragons

This is the latest (unstable) version of this documentation, which may document features not available in or compatible with released stable versions of Godot.

Checking the stable version of the documentation...

Використання обчислювальних шейдерів

Цей підручник проведе вас через процес створення мінімального обчислювального шейдера. Але спочатку трохи фону про обчислювальні шейдери та про те, як вони працюють із Godot.

Примітка

Цей посібник передбачає, що ви загалом знайомі з шейдерами. Якщо ви новачок у шейдерах, будь ласка, прочитайте Знайомство з шейдерами і ваш перший шейдер перед тим, як продовжити цей посібник.

Обчислювальний шейдер — це особливий тип шейдерної програми, яка орієнтована на програмування загального призначення. Іншими словами, вони більш гнучкі, ніж вершинні шейдери та фрагментні шейдери, оскільки вони не мають фіксованої мети (тобто трансформації вершин або запису кольорів у зображення). На відміну від фрагментних і вершинних шейдерів, у обчислювальних шейдерах дуже мало відбувається за лаштунками. Код, який ви пишете, — це те, що виконує GPU, і зовсім небагато іншого. Це може зробити їх дуже корисним інструментом для перевантаження важких обчислень на GPU.

Тепер давайте почнемо зі створення короткого обчислювального шейдера.

Спочатку у зовнішньому текстовому редакторі за вашим вибором створіть новий файл під назвою compute_example.glsl у папці вашого проекту. Коли ви пишете обчислювальні шейдери в Godot, ви пишете їх безпосередньо в GLSL. Мова шейдерів Godot базується на GLSL. Якщо ви знайомі зі звичайними шейдерами в Godot, наведений нижче синтаксис здасться вам дещо знайомим.

Примітка

Обчислювальні шейдери можна використовувати лише з рендерів на основі RenderingDevice (Forward+ або Mobile renderer). Щоб дотримуватися цього підручника, переконайтеся, що ви використовуєте Forward+ або Mobile рендерер. Налаштування для якого розташовано у верхньому правому куті редактора.

Зауважте, що підтримка обчислювальних шейдерів зазвичай погана на мобільних пристроях (через помилки драйверів), навіть якщо вони технічно підтримуються.

Погляньмо на цей код обчислювального шейдера:

#[compute]
#version 450

// Invocations in the (x, y, z) dimension
layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;

// A binding to the buffer we create in our script
layout(set = 0, binding = 0, std430) restrict buffer MyDataBuffer {
    float data[];
}
my_data_buffer;

// The code we want to execute in each invocation
void main() {
    // gl_GlobalInvocationID.x uniquely identifies this invocation across all work groups
    my_data_buffer.data[gl_GlobalInvocationID.x] *= 2.0;
}

Цей код приймає масив чисел з плаваючою точкою, множить кожен елемент на 2 і зберігає результати назад у буферному масиві. Тепер давайте розглянемо це рядок за рядком.

#[compute]
#version 450

Ці два рядки повідомляють про дві речі:

Наступний код є обчислювальним шейдером. Це підказка щодо Godot, яка потрібна редактору для правильного імпорту файлу шейдера.

Код використовує GLSL версії 450.

Вам ніколи не доведеться змінювати ці два рядки для ваших власних обчислювальних шейдерів.

// Invocations in the (x, y, z) dimension
layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;

Далі ми повідомляємо кількість викликів, які будуть використовуватися в кожній робочій групі. Виклики — це екземпляри шейдера, які виконуються в одній робочій групі. Коли ми запускаємо обчислювальний шейдер із ЦП, ми повідомляємо йому, скільки робочих груп потрібно запустити. Робочі групи працюють паралельно одна одній. Під час роботи однієї робочої групи ви не можете отримати доступ до інформації в іншій робочій групі. Однак виклики в одній робочій групі можуть мати обмежений доступ до інших викликів.

Подумайте про робочі групи та виклики як про гігантський вкладений цикл for.

for (int x = 0; x < workgroup_size_x; x++) {
  for (int y = 0; y < workgroup_size_y; y++) {
     for (int z = 0; z < workgroup_size_z; z++) {
        // Each workgroup runs independently and in parallel.
        for (int local_x = 0; local_x < invocation_size_x; local_x++) {
           for (int local_y = 0; local_y < invocation_size_y; local_y++) {
              for (int local_z = 0; local_z < invocation_size_z; local_z++) {
                 // Compute shader runs here.
              }
           }
        }
     }
  }
}

Робочі групи та виклики — це розширена тема. Наразі пам’ятайте, що ми будемо виконувати два виклики для кожної робочої групи.

// A binding to the buffer we create in our script
layout(set = 0, binding = 0, std430) restrict buffer MyDataBuffer {
    float data[];
}
my_data_buffer;

Тут ми надаємо інформацію про пам’ять, до якої матиме доступ обчислювальний шейдер. Властивість layout дозволяє нам вказати шейдеру, де шукати буфер, пізніше нам потрібно буде зіставити ці позиції set і binding з боку ЦП.

Ключове слово restrict повідомляє шейдеру, що доступ до цього буфера буде доступний лише з одного місця цього шейдера. Іншими словами, ми не будемо прив’язувати цей буфер до іншого індексу set або binding. Це важливо, оскільки дозволяє компілятору шейдерів оптимізувати код шейдера. Завжди використовуйте restrict, коли це можливо.

Це нерозмірний буфер, що означає, що він може бути будь-якого розміру. Тому нам потрібно бути обережними, щоб не читати з індексу, який перевищує розмір буфера.

// The code we want to execute in each invocation
void main() {
    // gl_GlobalInvocationID.x uniquely identifies this invocation across all work groups
    my_data_buffer.data[gl_GlobalInvocationID.x] *= 2.0;
}

Нарешті, ми пишемо функцію main, у якій і відбувається вся логіка. Ми отримуємо доступ до позиції в буфері зберігання за допомогою вбудованих змінних gl_GlobalInvocationID. gl_GlobalInvocationID дає глобальний унікальний ідентифікатор для поточного виклику.

Щоб продовжити, запишіть наведений вище код у щойно створений файл compute_example.glsl.

Створіть локальний RenderingDevice

Щоб взаємодіяти з обчислювальним шейдером і виконувати його, нам потрібен скрипт. Створіть новий скрипт вибраною вами мовою та прикріпіть його до будь-якого вузла вашої сцени.

Тепер для виконання нашого шейдера нам потрібен локальний class_RenderingDevice, який можна створити за допомогою class_RenderingServer:

# Create a local rendering device.
var rd := RenderingServer.create_local_rendering_device()

// Create a local rendering device.
var rd = RenderingServer.CreateLocalRenderingDevice();

Після цього ми можемо завантажити щойно створений файл шейдера compute_example.glsl і створити його попередньо скомпільовану версію за допомогою цього:

# Load GLSL shader
var shader_file := load("res://compute_example.glsl")
var shader_spirv: RDShaderSPIRV = shader_file.get_spirv()
var shader := rd.shader_create_from_spirv(shader_spirv)

// Load GLSL shader
var shaderFile = GD.Load<RDShaderFile>("res://compute_example.glsl");
var shaderBytecode = shaderFile.GetSpirV();
var shader = rd.ShaderCreateFromSpirV(shaderBytecode);

Попередження

Локальні RenderingDevices не можна налагодити за допомогою таких інструментів, як RenderDoc.

Надайте вхідні дані

Як ви пам’ятаєте, ми хочемо передати вхідний масив у наш шейдер, помножити кожен елемент на 2 і отримати результати.

Нам потрібно створити буфер для передачі значень у обчислювальний шейдер. Ми маємо справу з масивом float, тому для цього прикладу будемо використовувати буфер зберігання. Буфер зберігання займає масив байтів і дозволяє центральному процесору передавати дані в і з графічного процесора.

Тож давайте ініціалізуємо масив із плаваючими числами та створимо буфер зберігання:

# Prepare our data. We use floats in the shader, so we need 32 bit.
var input := PackedFloat32Array([1, 2, 3, 4, 5, 6, 7, 8, 9, 10])
var input_bytes := input.to_byte_array()

# Create a storage buffer that can hold our float values.
# Each float has 4 bytes (32 bit) so 10 x 4 = 40 bytes
var buffer := rd.storage_buffer_create(input_bytes.size(), input_bytes)

// Prepare our data. We use floats in the shader, so we need 32 bit.
float[] input = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10];
var inputBytes = new byte[input.Length * sizeof(float)];
Buffer.BlockCopy(input, 0, inputBytes, 0, inputBytes.Length);

// Create a storage buffer that can hold our float values.
// Each float has 4 bytes (32 bit) so 10 x 4 = 40 bytes
var buffer = rd.StorageBufferCreate((uint)inputBytes.Length, inputBytes);

З буфером на місці нам потрібно сказати пристрою рендерингу використовувати цей буфер. Для цього нам потрібно буде створити уніформу (як у звичайних шейдерах) і призначити її набору уніформ, який ми можемо передати нашому шейдеру пізніше.

# Create a uniform to assign the buffer to the rendering device
var uniform := RDUniform.new()
uniform.uniform_type = RenderingDevice.UNIFORM_TYPE_STORAGE_BUFFER
uniform.binding = 0 # this needs to match the "binding" in our shader file
uniform.add_id(buffer)
var uniform_set := rd.uniform_set_create([uniform], shader, 0) # the last parameter (the 0) needs to match the "set" in our shader file

// Create a uniform to assign the buffer to the rendering device
var uniform = new RDUniform
{
    UniformType = RenderingDevice.UniformType.StorageBuffer,
    Binding = 0
};
uniform.AddId(buffer);
var uniformSet = rd.UniformSetCreate([uniform], shader, 0);

Визначення обчислювального конвеєра

Наступним кроком є створення набору інструкцій, які може виконувати наш графічний процесор. Для цього нам потрібен конвеєр і список обчислень.

Щоб обчислити результат, нам потрібно виконати такі кроки:

Створіть новий конвеєр.
Почніть список інструкцій, які має виконати наш графічний процесор.
Прив’яжіть наш список обчислень до нашого конвеєра
Прив’яжіть нашу буферну форму до нашого трубопроводу
Укажіть, скільки робочих груп використовувати
Завершіть список інструкцій

# Create a compute pipeline
var pipeline := rd.compute_pipeline_create(shader)
var compute_list := rd.compute_list_begin()
rd.compute_list_bind_compute_pipeline(compute_list, pipeline)
rd.compute_list_bind_uniform_set(compute_list, uniform_set, 0)
rd.compute_list_dispatch(compute_list, 5, 1, 1)
rd.compute_list_end()

// Create a compute pipeline
var pipeline = rd.ComputePipelineCreate(shader);
var computeList = rd.ComputeListBegin();
rd.ComputeListBindComputePipeline(computeList, pipeline);
rd.ComputeListBindUniformSet(computeList, uniformSet, 0);
rd.ComputeListDispatch(computeList, xGroups: 5, yGroups: 1, zGroups: 1);
rd.ComputeListEnd();

Зауважте, що ми відправляємо обчислювальний шейдер із 5 робочими групами на осі X і по одній на інших. Оскільки у нас є 2 локальні виклики на осі X (зазначені в нашому шейдері), загалом буде запущено 10 викликів обчислювального шейдера. Якщо ви читаєте або записуєте індекси за межами діапазону вашого буфера, ви можете отримати доступ до пам’яті поза контролем ваших шейдерів або до частин інших змінних, що може спричинити проблеми з деяким обладнанням.

Виконайте обчислювальний шейдер

Після всього цього ми майже закінчили, але нам ще потрібно виконати наш конвеєр. Поки що ми записали лише те, що ми хотіли б, щоб GPU зробив; ми фактично не запускали шейдерну програму.

Щоб виконати наш обчислювальний шейдер, нам потрібно передати конвеєр до GPU та дочекатися завершення виконання:

# Submit to GPU and wait for sync
rd.submit()
rd.sync()

// Submit to GPU and wait for sync
rd.Submit();
rd.Sync();

В ідеалі ви б не викликали sync() для синхронізації RenderingDevice одразу, оскільки це змусить центральний процесор чекати завершення роботи графічного процесора. У нашому прикладі ми синхронізуємо відразу, тому що ми хочемо, щоб наші дані були доступні для читання відразу. Загалом, вам потрібно зачекати принаймні 2-3 кадри перед синхронізацією, щоб графічний процесор міг працювати паралельно з центральним процесором.

Попередження

Довгі обчислення можуть призвести до «збою» графічних драйверів Windows через TDR, який запускає Windows. Це механізм, який повторно ініціалізує графічний драйвер після того, як мине певний проміжок часу без жодної активності графічного драйвера (зазвичай від 5 до 10 секунд).

Залежно від тривалості виконання обчислювального шейдера, вам може знадобитися розділити його на кілька відправлень, щоб скоротити час, який займає кожне відправлення, і зменшити ймовірність запуску TDR. Враховуючи, що TDR залежить від часу, повільніші графічні процесори можуть бути більш схильні до TDR під час роботи певного обчислювального шейдера порівняно зі швидшим графічним процесором.

Отримання результатів

Можливо, ви помітили, що у прикладі шейдера ми змінили вміст буфера зберігання. Іншими словами, шейдер прочитав наш масив і знову зберіг дані в тому ж масиві, тому наші результати вже є. Давайте отримаємо дані та надрукуємо результати на нашу консоль.

# Read back the data from the buffer
var output_bytes := rd.buffer_get_data(buffer)
var output := output_bytes.to_float32_array()
print("Input: ", input)
print("Output: ", output)

// Read back the data from the buffers
var outputBytes = rd.BufferGetData(buffer);
var output = new float[input.Length];
Buffer.BlockCopy(outputBytes, 0, output, 0, outputBytes.Length);
GD.Print("Input: ", string.Join(", ", input));
GD.Print("Output: ", string.Join(", ", output));

Звільнення пам'яті

Змінні buffer, pipeline та uniform_set, які ми використовували, є змінними типу class_RID. Оскільки RenderingDevice задуманий як API нижчого рівня, RID не звільняються автоматично. Це означає, що після завершення використання buffer або будь-якого іншого RID ви відповідаєте за звільнення його пам'яті вручну за допомогою методу free_rid() класу RenderingDevice.

Таким чином, у вас є все необхідне, щоб розпочати роботу з обчислювальними шейдерами.

Дивись також

Репозиторій демонстраційних проектів містить Compute Shader Heightmap demo Цей проект виконує генерацію зображення висотної карти окремо на CPU та GPU, що дозволяє порівняти, як подібний алгоритм може бути реалізований двома різними способами (у більшості випадків реалізація на GPU є швидшою).