Заполнение MTLBuffer 16-битными числами

Я заполняю MTLBuffer векторами float2. Буфер создается и заполняется так:

struct Particle {
   var position: float2
   ...
}

let particleCount = 100000
let bufferSize = MemoryLayout<Particle>.stride * particleCount
particleBuffer = device.makeBuffer(length: bufferSize)!

var pointer = particleBuffer.contents().bindMemory(to: Particle.self, capacity: particleCount)
pointer = pointer.advanced(by: currentParticles)
pointer.pointee.position = [x, y]

В моем металлическом файле доступ к буферу осуществляется следующим образом:

struct Particle {
   float2 position;
   ...
};

kernel void compute(device Particle *particles [[buffer(0)]], … ) 

Мне нужно использовать поплавки с половинной точностью в моем вычислительном ядре Metal. Со стороны Metal это так же просто, как указать half2 для типа данных.

Что касается процессора, каков наилучший способ заполнения буфера с плавающими половинной точностью?

1 ответ

Решение

Поплавки половинной точности в Swift очень неудобны, так как нет Float16 типа еще (хотя один был предложен) и нестандартный __fp16 Тип, поддерживаемый Clang, также не полностью поддерживается в Swift.

Тем не менее, с помощью магии выкалывающих и соединяющих заголовки, вы сможете собрать работоспособное решение.

Основной подход заключается в следующем: в заголовке Objective C объявить half2 типа с двумя uint16_t члены. Это будет наш тип хранения. Также объявите функцию, которая принимает float и записывает его так, как если бы он был __fp16 на указатель наuint16_t:

typedef struct {
    uint16_t x, y;
} half2;

static void storeAsF16(float value, uint16_t *_Nonnull pointer) { *(__fp16 *)pointer = value; }

Вернувшись в Swift, вы можете объявить typealias и использовать его в своем определении структуры частиц:

typealias Half2 = half2

struct Particle {
    var position: Half2
}

(Здесь я набираю текст из строчных букв в имя Swiftier; вы можете пропустить это и просто назвать тип Obj-C Half2 если хочешь).

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

var pointer = particleBuffer.contents().bindMemory(to: Half2.self, capacity: particleCount)

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

var x: UInt16 = 0
var y: UInt16 = 0
storeAsF16(1.0, &x) // 0x3c00
storeAsF16(0.5, &y) // 0x3800

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

pointer.pointee = Half2(x: x, y: y)

Обратите внимание, что этот подход не является ни переносимым, ни безопасным, особенно потому, что Swift не дает никаких гарантий относительно структуры элементов struct. Могут быть и другие, менее громоздкие подходы; это то, что работало для меня в прошлом.

Другие вопросы по тегам