Металлический вычислительный трубопровод нелепо медленный - PullRequest
0 голосов
/ 27 ноября 2018

Я увидел возможность улучшить производительность своего приложения с помощью вычислительного конвейера Metal.Тем не менее, мое первоначальное тестирование показало, что конвейер вычислений был абсурдно медленным (по крайней мере, на старом устройстве).

Поэтому я сделал пример проекта для сравнения производительности вычислений и визуализации конвейеров.Программа берет исходную текстуру 2048 x 2048 и преобразует ее в оттенки серого в текстуре назначения.

На iPhone 5S для преобразования шейдера фрагмента потребовалось 3 мс.Тем не менее, вычислительному ядру потребовалось 177 мсек, чтобы сделать то же самое.Это в 59 раз больше !!!

enter image description here

Каков ваш опыт работы с конвейером вычислений на старом устройстве?Это не абсурдно медленно?

Вот мои фрагменты и функции вычисления:

// Grayscale Fragment Function
fragment half4 grayscaleFragment(RasterizerData in [[stage_in]],
                                 texture2d<half> inTexture [[texture(0)]])
{
    constexpr sampler textureSampler;

    half4 inColor  = inTexture.sample(textureSampler, in.textureCoordinate);
    half  gray     = dot(inColor.rgb, kRec709Luma);
    return half4(gray, gray, gray, 1.0);
}


// Grayscale Kernel Function
kernel void grayscaleKernel(uint2 gid [[thread_position_in_grid]],
                            texture2d<half, access::read> inTexture [[texture(0)]],
                            texture2d<half, access::write> outTexture [[texture(1)]])
{
    half4 inColor  = inTexture.read(gid);
    half  gray     = dot(inColor.rgb, kRec709Luma);
    outTexture.write(half4(gray, gray, gray, 1.0), gid);
}

Методы вычисления и рендеринга

- (void)compute {

    id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];

    // Compute encoder
    id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
    [computeEncoder setComputePipelineState:_computePipelineState];
    [computeEncoder setTexture:_srcTexture atIndex:0];
    [computeEncoder setTexture:_dstTexture atIndex:1];
    [computeEncoder dispatchThreadgroups:_threadgroupCount threadsPerThreadgroup:_threadgroupSize];
    [computeEncoder endEncoding];

    [commandBuffer commit];

    [commandBuffer waitUntilCompleted];
}


- (void)render {

    id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];

    // Render pass descriptor
    MTLRenderPassDescriptor *renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor];
    renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionDontCare;
    renderPassDescriptor.colorAttachments[0].texture = _dstTexture;
    renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore;

    // Render encoder
    id<MTLRenderCommandEncoder> renderEncoder = [commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
    [renderEncoder setRenderPipelineState:_renderPipelineState];
    [renderEncoder setFragmentTexture:_srcTexture atIndex:0];
    [renderEncoder drawPrimitives:MTLPrimitiveTypeTriangleStrip vertexStart:0 vertexCount:4];
    [renderEncoder endEncoding];

    [commandBuffer commit];

    [commandBuffer waitUntilCompleted];
}

И настройка металла:

- (void)setupMetal
{
    // Get metal device
    _device = MTLCreateSystemDefaultDevice();

    // Create the command queue
    _commandQueue = [_device newCommandQueue];

    id<MTLLibrary> defaultLibrary = [_device newDefaultLibrary];

    // Create compute pipeline state
    _computePipelineState = [_device newComputePipelineStateWithFunction:[defaultLibrary newFunctionWithName:@"grayscaleKernel"] error:nil];

    // Create render pipeline state
    MTLRenderPipelineDescriptor *pipelineStateDescriptor = [[MTLRenderPipelineDescriptor alloc] init];
    pipelineStateDescriptor.vertexFunction = [defaultLibrary newFunctionWithName:@"vertexShader"];
    pipelineStateDescriptor.fragmentFunction = [defaultLibrary newFunctionWithName:@"grayscaleFragment"];
    pipelineStateDescriptor.colorAttachments[0].pixelFormat = MTLPixelFormatBGRA8Unorm;
    _renderPipelineState = [_device newRenderPipelineStateWithDescriptor:pipelineStateDescriptor error:nil];

    // Create source and destination texture descriptor
    // Since the compute kernel function doesn't check if pixels are within the bounds of the destination texture, make sure texture width
    // and height are multiples of the pipeline threadExecutionWidth and (threadExecutionWidth / maxTotalThreadsPerThreadgroup) respectivly.
    MTLTextureDescriptor *textureDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatBGRA8Unorm
                                                                                                 width:2048
                                                                                                height:2048
                                                                                             mipmapped:NO];
    // Create source texture
    textureDescriptor.usage = MTLTextureUsageShaderRead;
    _srcTexture = [_device newTextureWithDescriptor:textureDescriptor];

    // Create description texture
    textureDescriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageRenderTarget;
    _dstTexture = [_device newTextureWithDescriptor:textureDescriptor];

    // Set the compute kernel's threadgroup size
    NSUInteger threadWidth = _computePipelineState.threadExecutionWidth;
    NSUInteger threadMax = _computePipelineState.maxTotalThreadsPerThreadgroup;
    _threadgroupSize = MTLSizeMake(threadWidth, threadMax / threadWidth, 1);

     // Set the compute kernel's threadgroup count
    _threadgroupCount.width  = (_srcTexture.width  + _threadgroupSize.width -  1) / _threadgroupSize.width;
    _threadgroupCount.height = (_srcTexture.height + _threadgroupSize.height - 1) / _threadgroupSize.height;
    _threadgroupCount.depth = 1;
}

1 Ответ

0 голосов
/ 28 ноября 2018

Металлический вычислительный конвейер непригоден для устройств CPU / GPU класса A7.Тот же конвейер вычислений имеет отличную производительность на A8 и более новых устройствах.Ваши варианты для решения этой проблемы состоят в том, чтобы создать фрагменты шейдеров для устройств A7 и использовать вычислительную логику для всех новых устройств, или вы можете экспортировать вычисления в ЦП на A7 (есть как минимум 2 ЦП с этим классом устройств).Вы также можете просто использовать все фрагментные шейдеры для всех устройств, но гораздо более высокая производительность сложного кода возможна с вычислительными ядрами, так что об этом стоит подумать.

...