Category Archives: Metal

Don’t reuse the same buffer to pass parameters inside a loop.

Posted on by
Reply

So here’s a mistake I made with the Metal API.

Suppose you have a loop where you’re constructing multiple encoders, one encoder per loop.

And you need to pass a parameter–say, an integer–into each encoder.

So you write the following:

id<MTLBuffer> buffer = [self.device newBufferWithLength:sizeof(uint16_t) options:MTLResourceOptionCPUCacheModeDefault];

for (uint16_t i = 0; i < 5; ++i) {
    id<MTLComputeCommandEncoder> compute = [buffer computeCommandEncoder];
    ... blah blah blah ...
    memmove(buffer.contents, &i, sizeof(i));
    [compute setBuffer:buffer offset:0 atIndex:MyKernelIndex];
    ... blah blah blah ...
    [compute dispatchThread...];
    [compute endEncoding];
}

If you run this, I discovered that all five invocations of the kernel will result in the two-byte value at MyKernelIndex to be set to 4–the last value seen in i as we loop.

Why?

Because the same buffer is reused across all five invocations, and because the Metal code isn’t executed until after the entire buffer is committed–the last value passed in is the value that will be used across all invocations.

But if this is replaced with:

for (uint16_t i = 0; i < 5; ++i) {
    id<MTLComputeCommandEncoder> compute = [buffer computeCommandEncoder];
    ... blah blah blah ...
    [compute setBytes:&i length:sizeof(i) atIndex:MyKernelIndex];
    ... blah blah blah ...
    [compute dispatchThread...];
    [compute endEncoding];
}

Each invocation gets a unique value for i.

Just something to watch out for in a deferred execution model.

Updated “Metal Introduction” Document.

Posted on by
Reply

I’ve added a section on showing the processes in Metal for implementing Constructive Solid Geometry on the fly using the algorithms outlined in the paper An improved z-buffer CSG rendering algorithm, with example code uploaded to GitHub.

My goal is to eventually turn this into a CSG library for Metal.

The updated paper (with the additional section) can be downloaded from here: Metal: An Introduction, and updates the paper from my prior post.

And when you put it all together you should get:

ScreenShot2

Learning the Metal API

Posted on by
Reply

So I’m in the process of learning the Metal API, which is Apple’s replacement for OpenGL ES. The principles are fairly similar, though the Metal API is much lower level.

There are several web sites devoted to Metal, but my eventual goal is to implement image-based CSG (Constructive Solid Geometry) in Metal for an update to the Kythera application.

And that requires a deeper understanding of Metal than most introductions which seem to stop at drawing an object on the screen and perhaps adding a texture map to the object.

The best way to learn something is to try to explain it–so I’ve started writing a document showing how to build a Macintosh-based Metal API application in Objective C, and going from a blank screen to a deferred shading example.

In this case, the deferred shading example results in a rotating teapot with fairy lights and indirect illumination rendering at 60 frames/second:

Teapot Rendering Example

The sample code is uploaded at GitHub, with the different examples in their own branches.

And the first draft of the Metal Introduction Document (as a PDF with links to relevant documents) can be downloaded from here.

Feedback is appreciated.