CODEWITHSUNDEEP
swiftswift4metalhalf-precision-float
Jeshua Lacock
Jeshua Lacock

Reputation: 6658

Populating MTLBuffer with 16-bit Floats

I am populating an MTLBuffer with float2 vectors. The buffer is being created and populated like this:

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]

In my Metal file the buffer is being accessed like this:

struct Particle {
   float2 position;
   ...
};

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

I need to use half precision floats in my Metal compute kernel. On the Metal side, it is as easy as specifying half2 for the data type.

On the CPU side, what is the best way to populate a buffer with half precision floats?

Upvotes: 4

Views: 1046

Answers (1)

warrenm
warrenm

Reputation: 31782

Half-precision floats in Swift are very awkward, since there is no Float16 type yet (though one has been proposed) and the nonstandard __fp16 type supported by Clang isn't fully supported in Swift either.

Through the magic of type-punning and bridging headers, however, you may be able to cobble together a workable solution.

The basic approach is this: In an Objective-C header, declare a half2 type with two uint16_t members. These will be our storage type. Also declare a function that takes a float and writes it as if it were a __fp16 to a pointer-to-uint16_t:

typedef struct {
    uint16_t x, y;
} half2;

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

Back in Swift, you can declare a typealias and use it in your particle struct definition:

typealias Half2 = half2

struct Particle {
    var position: Half2
}

(Here I'm typealiasing from the lower-case type to a Swiftier name; you could skip this and just name the Obj-C type Half2 if you prefer).

Instead of binding to the particle type, you'll need to bind the buffer to your half vector type instead:

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

When we use our utility function to store a float, the bit pattern for the corresponding half value gets written to the UInt16:

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

Now that we have correctly-formated half values in this pair of variables, we can write them into the buffer:

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

Note that this approach is neither portable nor safe, especially because Swift doesn't make any guarantees about struct member layout. There may be other less cumbersome approaches, too; this is just what has worked for me in the past.

Upvotes: 4

Related Questions