Monday, August 20, 2018

Performance of metal function multiple call

i make rigid body simulation for iPhone/iPad with using Apple Metal. To do this, i need to make many calls of kernel functions, and i see, that it takes a long time, opposite to CUDA for example. I implemented Metal kernel function call, like it describes in Apple tutorial

let commandQueue = device.newCommandQueue()

var commandBuffers:[MTLCommandBuffer]=[]
var gpuPrograms:[MTLFunction]=[]
var computePipelineFilters:[MTLComputePipelineState]=[]
var computeCommandEncoders:[MTLComputeCommandEncoder]=[]

//here i fill all arrays for my command queue
//and next i execute it 

let threadsPerGroup = MTLSize(width:1,height:1,depth:1)
let numThreadgroups = MTLSize(width:threadsAmount, height:1, depth:1)

for computeCommandEncoder in computeCommandEncoders
{
    computeCommandEncoder.dispatchThreadgroups(numThreadgroups, threadsPerThreadgroup: threadsPerGroup)
}

for computeCommandEncoder in computeCommandEncoders
{
    computeCommandEncoder.endEncoding()
}

for commandBuffer in commandBuffers
{
    commandBuffer.enqueue()
}

for commandBuffer in commandBuffers
{
    commandBuffer.commit()
}

for commandBuffer in commandBuffers
{
    commandBuffer.waitUntilCompleted()
}

I am do up to few dozens metal kernel functions every frame, and it works too slow. I tested it with empty kernel functions - and it shows me, that the problem are in Swift part of execution. I mean, when i want to execute kernel function in CUDA, i just call it like usual function and it works very fast. But here i must make many actions for every execution of every function every frame. May be i don't know something, but i want create all additional objects one time, and then just make something like

 commandQueue.execute()

to execute all kernel functions.

Am i rights in my actions to execute many kernel functions, or there is some other way to do it faster?

Solved

I have a few projects that use multiple shaders in a single step. I only create a single buffer and encoder but multiple pipeline states; one for each compute function.

Remember that MTLCommandQueue is persistent, so only needs to be created once, so my MetalKit View's drawRect() function is roughly this (there are more shaders and textures being passed between them, but you get an idea of structure):

let commandBuffer = commandQueue.commandBuffer()
let commandEncoder = commandBuffer.computeCommandEncoder()

commandEncoder.setComputePipelineState(advect_pipelineState)
commandEncoder.dispatchThreadgroups(threadgroupsPerGrid, 
    threadsPerThreadgroup: threadsPerThreadgroup)

commandEncoder.setComputePipelineState(divergence_pipelineState)
commandEncoder.dispatchThreadgroups(threadgroupsPerGrid, 
    threadsPerThreadgroup: threadsPerThreadgroup)

[...]

commandEncoder.endEncoding()
commandBuffer.commit()

My code actually iterates over one of the shaders twenty times and still runs pretty nippily, so if you reorganise and to follow this structure with a single buffer and a single encoder and only call endEncoding() and commit() once per pass, you may see an improvement in performance.

May being the operative word :)


No comments:

Post a Comment