Вызов второго ядра Metal в кадре рендеринга происходит медленно

У меня есть анимация, которая вызывает металлические ядра, чтобы вычислить некоторые данные для меня.

Если я назову только одно простое ядро ​​avg. время вызова ядра составляет 4 мс. Но когда я вызываю второе ядро ​​после первого (очень похожее, простое ядро), второй вызов ядра составляет 40 мс или более … Это что-то в корне не так с моим кодом?

Я настраивал устройство, очередь, библиотеку один раз перед запуском анимации. Я также создаю состояние конвейера для каждого ядра впереди.

// globals created once in the beginning
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
id<MTLCommandQueue> queue = [device newCommandQueueWithMaxCommandBufferCount:1];

MTLCompileOptions *options = [MTLCompileOptions new];
options.fastMathEnabled = YES;
id<MTLLibrary> metal_library = [device newLibraryWithSource:@(MetalKernel_code) options:options error:&err]
[options release];

id<MTLFunction> kernel1 = [[metal_library newFunctionWithName:[NSString stringWithUTF8String:@"kernel1"]/* constantValues : constantValues */] autorelease]);

id<MTLComputePipelineState> pipeline_state_1 = [device newComputePipelineStateWithFunction:kernel1 error:&err];

id<MTLFunction> kernel2 = [[metal_library newFunctionWithName:[NSString stringWithUTF8String:@"kernel2"]/* constantValues : constantValues */] autorelease];
id<MTLComputePipelineState> pipeline_state_2 = [device newComputePipelineStateWithFunction:kernel2 error:&err];

И теперь каждый кадр, который я называю ядром, выглядит так:

// invoking first kernel:

id<MTLBuffer> srcDeviceBuf = reinterpret_cast<id<MTLBuffer> >(const_cast<float *>(src));
id<MTLBuffer> dstDeviceBuf = reinterpret_cast<id<MTLBuffer> >(dst);

Kernel1Params params{w, h};
id<MTLBuffer> paramBuf = [[device newBufferWithBytes:&params length:sizeof(Kernel1Params) options:MTLResourceStorageModeManaged]
autorelease];

id<MTLCommandBuffer> commandBuffer = [queue commandBuffer];
commandBuffer.label = [NSString stringWithFormat:@"kernel1"];

id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
[computeEncoder setComputePipelineState:pipeline_state_1];

auto exeWidth = [pipeline_state_1 threadExecutionWidth];
MTLSize threadGroupCount = MTLSizeMake(exeWidth, 1, 1);
MTLSize threadGroups = MTLSizeMake((w + exeWidth - 1) / exeWidth, static_cast<uint32_t>(h), 1);

[computeEncoder setBuffer:srcDeviceBuf offset:0 atIndex:0];
[computeEncoder setBuffer:dstDeviceBuf offset:0 atIndex:1];
[computeEncoder setBuffer:paramBuf offset:0 atIndex:2];

[computeEncoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupCount];
[computeEncoder endEncoding];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];// invoking second kernel

Kernel2Params params{w_small, h_small, w_big, h_big};
id<MTLBuffer> srcDeviceBuf = reinterpret_cast<id<MTLBuffer> >(const_cast<float *>(src));
id<MTLBuffer> dstDeviceBuf = reinterpret_cast<id<MTLBuffer> >(const_cast<float *>(dst));
id<MTLBuffer> paramBuf = [[device newBufferWithBytes:&params length:sizeof(Kernel2Params) options:MTLResourceStorageModeManaged] autorelease];

id<MTLCommandBuffer> commandBuffer = [queue commandBuffer];
commandBuffer.label = [NSString stringWithFormat:@"Kernel2"];

id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
[computeEncoder setComputePipelineState:pipeline_state_2];

auto exeWidth = [pipeline_state_2 threadExecutionWidth];
MTLSize threadGroupCount = MTLSizeMake(exeWidth, 1, 1);
MTLSize threadGroups = MTLSizeMake((w_small + exeWidth - 1) / exeWidth, static_cast<uint32_t>(h_small), 1);

[computeEncoder setBuffer:srcDeviceBuf offset:0 atIndex:0];
[computeEncoder setBuffer:dstDeviceBuf offset:0 atIndex:1];
[computeEncoder setBuffer:paramBuf offset:0 atIndex:2];

[computeEncoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupCount];
[computeEncoder endEncoding];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];

// copy from GPU to CPU
NSData *outdata = [NSData dataWithBytesNoCopy:[dstDeviceBuf contents] length:(size_t) (w_small * h_small * 3) freeWhenDone:false];

[outdata getBytes:dst length:(size_t) (w_small * h_small * 3)];

Оба ядра могут быть такими же простыми, как копирование входного буфера в выходной буфер (1) и простое изменение размера изображения (2):

// kernel 1
kernel void Kernel1( device float const* src, device float* dst, device Kernel1Params *inValues, uint2 inXY [[ thread_position_in_grid ]])
{
if ((inXY.x >= w) || (inXY.y >= h)) {
return;
}
const int index = ((inXY.y * inValues->w) + inXY.x) * 4;

dst[index + 0] = src[index + 0];
dst[index + 1] = src[index + 1];
dst[index + 2] = src[index + 2];
dst[index + 3] = src[index + 3];
}

// kernel 2
kernel void Kernel2( device float const* src, device unsigned char* dst, device Kernel2Params *inValues, uint2 inXY [[ thread_position_in_grid ]])
{
if ((inXY.x >= w_small) || (inXY.y >= h_small)) {
return;
}
const int index = ((inXY.y * inValues->w_small) + inXY.x) * 3;
int index2;

if(inValues->w_small == inValues->w_big && inValues->h_small == inValues->h_big) {
index2 = ((inXY.y * inValues->w_big) + inXY.x) * 4;
} else {
const unsigned int x_big = (int)((float)(inXY.x * inValues->w_big) / (float)inValues->w_small);
const unsigned int y_big = inValues->h_big - (int)((float)(inXY.y * inValues->h_big) / (float)inValues->h_small) - 1;

index2 = (y_big * inValues->w_big + x_big) * 4;
}

dst[index + 0] = (unsigned char)(src[index2 + 2] * 255.0f);
dst[index + 1] = (unsigned char)(src[index2 + 1] * 255.0f);
dst[index + 2] = (unsigned char)(src[index2 + 0] * 255.0f);
}