Передача параметров в Metal Compute Kernel с использованием Swift 4

На стороне процессора у меня есть структура, которую я хочу передать вычислительному ядру:

  private struct BoundingBoxParameters {
    var x: Float = 0
    var y: Float = 0
    var width: Float = 0
    var height: Float = 0
    var levelOfDetail: Float = 1.0
    var dummy: Float = 1.0  // Needed for success
  }

Перед запуском ядра я передаю данные в MTLComputeCommandEncoder:

Вариант 1 (напрямую):

commandEncoder!.setBytes(&params, length: MemoryLayout<BoundingBoxParameters>.size, index: 0)

Вариант 2 (косвенно через MTLBuffer):

boundingBoxBuffer.contents().copyBytes(from: &params, count: MemoryLayout<BoundingBoxParameters>.size)
commandEncoder!.setBuffer(boundingBoxBuffer, offset: 0, index: 0)

Любой вариант работает нормально, если в структуре существует «фиктивная» переменная, но не работает, если «фиктивная» переменная не существует. Код не работает при вызове:

commandEncoder!.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount)

С ошибкой:

validateComputeFunctionArguments:820: failed assertion `Compute Function(resizeImage): argument params[0] from buffer(0) with offset(0) and length(20) has space for 20 bytes, but argument has a length(24).'

На стороне металлического ядра вот соответствующие фрагменты кода:

struct BoundingBoxParameters {
  float2 topLeft;
  float2 size;
  float levelOfDetail;
};

kernel void resizeImage(constant BoundingBoxParameters *params [[buffer(0)]],
                        texture2d<half, access::sample> sourceTexture [[texture(0)]],
                        texture2d<half, access::write> destTexture [[texture(1)]],
                        sampler samp [[sampler(0)]],
                        uint2 gridPosition [[thread_position_in_grid]]) {
  float2 destSize = float2(destTexture.get_width(0), destTexture.get_height(0));
  float2 sourceCoords = float2(gridPosition) / destSize;
  sourceCoords *= params->size;
  sourceCoords += params->topLeft;
  float lod = params->levelOfDetail;
  half4 color = sourceTexture.sample(samp, sourceCoords, level(lod));
  destTexture.write(color, gridPosition);
}

Я также получаю аналогичную проблему при попытке передать матрицу 3x3 другому вычислительному ядру. Он жалуется, что предоставлено 36 байтов, но ожидается 48.

У кого-нибудь есть идеи по этому вопросу?


kernel swift4 metalkit image-processing metal
person Labeno    schedule 10.02.2018    source источник

Ответы (1)


arrow_upward
8
arrow_downward

Прежде всего, я хочу отметить, что вам не следует использовать size когда вам нужно получить фактическую длину типа Swift, размещенного в памяти. Для этого следует использовать stride. Согласно Type Layout от Swift :

Окончательный размер и выравнивание — это размер и выравнивание агрегата. Шаг – это окончательный размер, округленный до выравнивания.

В этом ответе подробно рассказывается о расположении памяти в Swift, если вы хотите лучше понять тему.

Проблема в том, что ваш Metal struct, который использует float2, и Swift struct, который заменяет его двумя отдельными полями Float, имеют разные схемы памяти.

Размер (шаг в случае Swift) структуры должен быть кратен наибольшему выравниванию любого члена структуры. Наибольшее выравнивание в вашем Metal struct составляет 8 байт (выравнивание float2), поэтому в конце структуры после значения float есть дополнение.

struct BoundingBoxParameters {
    float2 topLeft; // 8 bytes
    float2 size; // 8 bytes
    float levelOfDetail; // 4 bytes
    // 4 bytes of padding so that size of struct is multiple 
    // of the largest alignment (which is 8 bytes)

}; // 24 bytes in total

Таким образом, ваш Metal struct на самом деле занимает 24 байта, как предполагает ошибка.

В то же время вашему Swift struct, имеющему наибольшее выравнивание в 4 байта, нужно всего 20 байтов.

private struct BoundingBoxParameters {
    var x: Float = 0 // 4 bytes
    var y: Float = 0 // 4 bytes
    var width: Float = 0 // 4 bytes
    var height: Float = 0 // 4 bytes
    var levelOfDetail: Float = 1.0 // 4 bytes
    // no need for any padding 

} // 20 bytes in total

Вот почему они оказываются несовместимыми друг с другом и полем dummy, компенсирующим 4 недостающих байта до Swift struct.

Чтобы решить эту проблему, я предлагаю вам использовать float2 из simd в Swift вместо Floats:

import simd 

private struct BoundingBoxParameters {
    var topLeft = float2(x: 0, y: 0)
    var size = float2(x: 0, y: 0)
    var levelOfDetail: Float = 1.0 
}

Не забудьте использовать MemoryLayout<BoundingBoxParameters>.stride (24 байта), чтобы получить длину вместо size (20 байтов).

То же самое касается случая матрицы 3x3: float3x3 в Metal имеет размер 48 байт и выравнивание 16 байт. Как я предполагаю, вы создали Swift struct с 9 Floats, которые будут иметь шаг/размер 36 байтов и выравнивание 4 байта. Отсюда и несоосность. Используйте matrix_float3x3 из simd.

В общем, для любых случаев, когда вы используете векторы или матрицы в Metal, вы должны использовать соответствующие типы simd в Swift.

person Lësha Turkowski    schedule 11.02.2018
comment
Спасибо за супер полезный ответ. Это решило проблему. - person Labeno; 11.02.2018
comment
Хотел бы я найти ваше предложение simd месяц назад. Очень хорошо объяснили и помогли. Спасибо! :) - person Kaolin Fire; 16.12.2018