Применение вычислительных шейдеров

В этом руководстве мы покажем вам процесс создания минимального вычислительного шейдера. Но сначала немного информации о вычислительных шейдерах и их работе с Godot.

Примечание

В этом руководстве предполагается, что вы в целом знакомы с шейдерами. Если вы новичок в этом деле, пожалуйста, прочтите Введение в шейдеры и your first shader, прежде чем продолжить изучение этого руководства.

Сompute (Вычислительный) шейдер — это особый тип шейдерной программы, ориентированный на программирование общего назначения. Другими словами, они более гибкие, чем vertex и fragment шейдеры, поскольку у них нет фиксированной цели (например, преобразования вершин или записи цветов в изображение). В отличие от vertex и fragment шейдеров, вычислительные шейдеры практически не выполняют никаких внутренних операций. Графический процессор выполняет только написанный вами код, и практически ничего больше. Это делает их очень полезным инструментом для переноса тяжёлых вычислений на графический процессор.

Теперь приступим к созданию короткого compute шейдера.

Сначала во внешнем текстовом редакторе по вашему выбору создайте новый файл compute_example.glsl в папке вашего проекта. Вычислительные шейдеры в Godot пишутся непосредственно на GLSL. Язык шейдеров Godot основан на GLSL. Если вы знакомы с обычными шейдерами в Godot, синтаксис ниже покажется вам знакомым.

Примечание

Вычислительные шейдеры можно использовать только из рендереров на базе RenderingDevice (рендер Forward+ или Mobile). Для выполнения этого руководства убедитесь, что вы используете рендер 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;

Затем мы сообщаем количество вызовов, которые будут использоваться в каждой рабочей группе. Вызовы — это экземпляры шейдера, работающие в одной рабочей группе. При запуске вычислительного шейдера из CPU мы сообщаем ему, сколько рабочих групп нужно запустить. Рабочие группы работают параллельно друг другу. При запуске одной рабочей группы доступ к информации в другой рабочей группе невозможен. Однако вызовы в одной рабочей группе могут иметь ограниченный доступ к другим вызовам.

Подумайте о рабочих группах и вызовах как о гигантском вложенном цикле 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

Для взаимодействия с вычислительным шейдером и его выполнения нам понадобится скрипт. Создайте новый скрипт на выбранном вами языке и прикрепите его к любому Узлу в вашей сцене.

Теперь для выполнения нашего шейдера нам нужен локальный RenderingDevice, который можно создать с помощью 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 и получить результат.

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

Итак, давайте инициализируем массив чисел с плавающей точкой и создадим буфер хранения:

# 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);

После того, как буфер готов, нам нужно указать устройству рендеринга использовать его. Для этого нам потребуется создать uniform (как в обычных шейдерах) и назначить его набору uniform, который мы позже передадим в наш шейдер.

# 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);

Определение вычислительного pipeline

Следующий шаг — создание набора инструкций, которые может выполнить наш графический процессор. Для этого нам понадобится конвейер и список вычислений.

Шаги, которые нам необходимо выполнить для вычисления результата:

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

# 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();

Обратите внимание, что мы запускаем вычислительный шейдер с пятью рабочими группами по оси X и одной по остальным. Поскольку у нас есть два локальных вызова по оси X (указанных в нашем шейдере), всего будет запущено 10 вызовов вычислительного шейдера. При чтении или записи по индексам за пределами диапазона буфера возможен доступ к памяти за пределами контроля шейдера или к частям других переменных, что может вызвать проблемы на некоторых устройствах.

Выполнить вычислительный шейдер

После всего этого мы почти закончили, но нам ещё нужно запустить наш конвейер. Пока что мы только записали, что именно должен делать графический процессор; мы ещё не запустили шейдерную программу.

Чтобы выполнить наш вычислительный шейдер, нам необходимо передать конвейер графическому процессору и дождаться завершения выполнения:

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

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

В идеале не следует вызывать sync() для синхронизации RenderingDevice сразу, так как это заставит CPU ждать завершения работы GPU. В нашем примере мы синхронизируемся сразу, потому что хотим, чтобы данные были доступны для чтения немедленно. В общем случае, перед синхронизацией следует подождать как минимум 2-3 кадра, чтобы GPU мог работать параллельно с CPU.

Предупреждение

Длительные вычисления могут привести к "crash" графических драйверов 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 представляют собой RID. Поскольку RenderingDevice задуман как низкоуровневый API, RID не освобождаются автоматически. Это означает, что после завершения использования buffer или любого другого RID вам необходимо вручную освободить память с помощью метода free_rid() класса RenderingDevice.

Теперь у вас есть все необходимое для начала работы с вычислительными шейдерами.

См. также

Репозиторий демонстрационных проектов содержит демонстрацию Compute Shader Heightmap Этот проект выполняет генерацию изображения карты высот отдельно на центральном процессоре и графическом процессоре, что позволяет сравнить, как один и тот же алгоритм может быть реализован двумя различными способами (при этом реализация на графическом процессоре в большинстве случаев оказывается быстрее).