Running compute kernel on portion of MTLBuffer?

Question

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)]], 
                    uint id [[thread_position_in_grid]] … )

I need to be able to compute a given range of the MTLBuffer. For example, is it possible to run the compute kernel on say starting from the 50,000 value and ending at 75,000 value?

It seems like the offset parameter allows me to specify the start position, but it does not have a length parameter.

I see there is this call:

setBuffers(_:offsets:range:)

Does the range specify what portion of the buffer to run? It seems like the range specifies what buffers are used and not the range of values to use.

Ken Thomases Ken Thomases · Accepted Answer · 2019-06-16T00:06:21

A compute shader doesn't run "on" a buffer (or portion). It runs on a grid, which is an abstract concept. As far as Metal is concerned, the grid isn't related to a buffer or anything else.

A buffer may be an input that your compute shader uses, but how it uses it is up to you. Metal doesn't know or care.

Here's my answer to a similar question.

So, the dispatch command you encode using a compute command encoder governs how many times your shader is invoked. It also dictates what thread_position_in_grid (and related) values each invocation receives. If your shader correlates each grid position to an element of an array backed by a buffer, then the number of threads you specify in your dispatch command governs how much of the buffer you end up accessing. (Again, this is not something Metal dictates; it's implicit in the way you code your shader.)

Now, to start at the 50,000th element, using an offset on the buffer binding to make that the effective start of the buffer from the point of view of the shader is a good approach. But it would also work to just add 50,000 to the index the shader computes when it accesses the buffer. And, if you only want to process 25,000 elements (75,000 minus 50,000), then just dispatch 25,000 threads.

Running compute kernel on portion of MTLBuffer?

1 Answers