передача изменяемого MTLBuffer между двумя быстродействующими кодировщиками металла - PullRequest
Новая
       10

передача изменяемого MTLBuffer между двумя быстродействующими кодировщиками металла

0 голосов
/ 09 июня 2018

В некоторых объективных классах c у меня есть буфер, который используется двумя последовательными металлическими конвейерами.Одно ядро ​​обрабатывает значения, основанные на яркости rgb изображения, и записывает их в изменяемый буфер, а следующее строит график формы сигнала на основе данных первого.Он основан на шейдере osx, описанном здесь ИСПОЛЬЗОВАНИЕ МЕТАЛЛИЧЕСКИХ ШЕЙДЕРОВ ПРОИЗВОДИТЕЛЬНОСТИ С ОСНОВНЫМ ИЗОБРАЖЕНИЕМ .

Мне удалось преобразовать это из OSX в IOS, который работает отлично, но сейчас я пытаюсь преобразовать его в быстрый, чтобы сделать его более гибким для меня в будущем.У меня есть другие фильтры, которые я хотел бы написать, и процесс перевода также научил меня тому, как работают металлические трубопроводы.

Вот где он ломается.На графике, который он рисует, одна горизонтальная цветная линия развевается вверх и вниз.Близко, но не сигара.Я задавался вопросом, было ли это как-то связано с тем, как были созданы буферы в быстром против цели c.Вот пример оригинала: 

size_t columnBufSize = sizeof(UInt64)*inputTexture.width*inputTexture.height;
id<MTLBuffer> columnDataRed = [kDevice newBufferWithLength:columnBufSize options:0];

, но когда я перевожу это по-быстрому, он работает не так, как ожидалось. 

 let columnBufSize: size_t = MemoryLayout<UInt64>.size * inputTexture!.width  * inputTexture!.height
 let columnDataRed: MTLBuffer  = kDevice.makeBuffer( length: columnBufSize, options: .storageModeShared)!

Похоже, что во второе передается только последнее значениеcomputeEncoder.

Что мне не хватает?

для завершения, вот полный код цели c 

+ (BOOL)processWithInputs:(NSArray<id<CIImageProcessorInput>> *)inputs arguments:(NSDictionary<NSString *,id> *)arguments output:(id<CIImageProcessorOutput>)output error:(NSError * _Nullable *)error
{
    id<MTLComputePipelineState> renderComputerState = kParadeComputePipelineState;



    id<CIImageProcessorInput> input = inputs.firstObject;

    id<MTLCommandBuffer> commandBuffer = output.metalCommandBuffer;
    commandBuffer.label = @"com.martinhering.WaveformKernel";
    id<MTLTexture> inputTexture = input.metalTexture;
    id<MTLTexture> outputTexture = output.metalTexture;

    MTLSize threadgroupCount = MTLSizeMake(inputTexture.width, inputTexture.height, 1);
    MTLSize _threadgroupSize = MTLSizeMake(16, 16, 1);

    threadgroupCount.width  = (inputTexture.width  + _threadgroupSize.width -  1) / _threadgroupSize.width;
    threadgroupCount.height = (inputTexture.height + _threadgroupSize.height - 1) / _threadgroupSize.height;


    size_t columnBufSize = sizeof(UInt64)*inputTexture.width*inputTexture.height;
    id<MTLBuffer> columnDataRed = [kDevice newBufferWithLength:columnBufSize options:0];
    id<MTLBuffer> columnDataGreen = [kDevice newBufferWithLength:columnBufSize options:0];
    id<MTLBuffer> columnDataBlue = [kDevice newBufferWithLength:columnBufSize options:0];

    id<MTLComputeCommandEncoder> computeEncoder;

    computeEncoder = [commandBuffer computeCommandEncoder];
    [computeEncoder setComputePipelineState:kWaveformComputePipelineState];
    [computeEncoder setTexture:inputTexture atIndex:0];
    [computeEncoder setBuffer:columnDataRed offset:0 atIndex:0];
    [computeEncoder setBuffer:columnDataGreen offset:0 atIndex:1];
    [computeEncoder setBuffer:columnDataBlue offset:0 atIndex:2];
    [computeEncoder setSamplerState:kSamplerState atIndex:0];
    [computeEncoder dispatchThreadgroups:threadgroupCount
                   threadsPerThreadgroup:_threadgroupSize];

    [computeEncoder endEncoding];


    computeEncoder = [commandBuffer computeCommandEncoder];
    [computeEncoder setComputePipelineState:renderComputerState];
    [computeEncoder setTexture:inputTexture atIndex:0];
    [computeEncoder setTexture:outputTexture atIndex:1];
    [computeEncoder setBuffer:columnDataRed offset:0 atIndex:0];
    [computeEncoder setBuffer:columnDataGreen offset:0 atIndex:1];
    [computeEncoder setBuffer:columnDataBlue offset:0 atIndex:2];
    [computeEncoder setSamplerState:kSamplerState atIndex:0];
    [computeEncoder dispatchThreadgroups:threadgroupCount
                   threadsPerThreadgroup:_threadgroupSize];

    [computeEncoder endEncoding];
    return YES;
}

Вот мой перевод 

override class func process(with inputs: [CIImageProcessorInput]?, arguments: [String : Any]?, output: CIImageProcessorOutput) throws {

    guard
        let kDevice = device,
        let commandBuffer = output.metalCommandBuffer,
        let input = inputs?.first,
        let defaultLibrary: MTLLibrary = kDevice.makeDefaultLibrary()
    else  {
        return
    }

     let samplerDescriptor = MTLSamplerDescriptor()
    let  kSamplerState = kDevice.makeSamplerState(descriptor: samplerDescriptor)

    samplerDescriptor.sAddressMode = .clampToEdge
    samplerDescriptor.tAddressMode = .clampToEdge
    samplerDescriptor.minFilter = .nearest
    samplerDescriptor.magFilter = .nearest
    samplerDescriptor.normalizedCoordinates = false


      var kWaveformComputePipelineState: MTLComputePipelineState?         
      var kParadeComputePipelineState: MTLComputePipelineState?


    if let aFunction = defaultLibrary.makeFunction(name: "scope_waveform_compute") {
        kWaveformComputePipelineState = try? kDevice.makeComputePipelineState(function: aFunction)
    }

    if let aFunction = defaultLibrary.makeFunction(name: "scope_waveform_parade") {
        kParadeComputePipelineState = try? kDevice.makeComputePipelineState(function: aFunction)
    }






    commandBuffer.label = "com.martinhering.WaveformKernel"

    weak var inputTexture: MTLTexture? = input.metalTexture
    weak var outputTexture: MTLTexture? = output.metalTexture

    var threadgroupCount: MTLSize = MTLSizeMake(inputTexture!.width, inputTexture!.height, 1)
    let threadgroupSize: MTLSize = MTLSizeMake(16, 16, 1)
    threadgroupCount.width = (inputTexture!.width  + threadgroupSize.width - 1) / threadgroupSize.width
    threadgroupCount.height = (inputTexture!.height + threadgroupSize.height - 1) / threadgroupSize.height


    let columnBufSize: size_t = MemoryLayout<UInt64>.size * inputTexture!.width  * inputTexture!.height

    let columnDataRed: MTLBuffer  = kDevice.makeBuffer( length: columnBufSize, options: .storageModeShared)!
    let columnDataGreen: MTLBuffer  = kDevice.makeBuffer( length: columnBufSize, options: .storageModeShared)!
    let columnDataBlue: MTLBuffer  = kDevice.makeBuffer( length: columnBufSize, options: .storageModeShared)!


    weak var computeEncoder: MTLComputeCommandEncoder?

    computeEncoder = commandBuffer.makeComputeCommandEncoder()
    computeEncoder?.setComputePipelineState(kWaveformComputePipelineState!)
    computeEncoder?.setTexture(inputTexture, index: 0)
    computeEncoder?.setBuffer(columnDataRed, offset: 0, index: 0)
    computeEncoder?.setBuffer(columnDataGreen, offset: 0, index: 1)
    computeEncoder?.setBuffer(columnDataBlue, offset: 0, index: 2)
    computeEncoder?.setSamplerState(kSamplerState, index: 0)
    computeEncoder?.dispatchThreadgroups(threadgroupCount, threadsPerThreadgroup: threadgroupSize)
    computeEncoder?.endEncoding()


    computeEncoder = commandBuffer.makeComputeCommandEncoder()
    computeEncoder?.setComputePipelineState(kParadeComputePipelineState!)
    computeEncoder?.setTexture(inputTexture, index: 0)
    computeEncoder?.setTexture(outputTexture, index: 1)
    computeEncoder?.setBuffer(columnDataRed, offset: 0, index: 0)
    computeEncoder?.setBuffer(columnDataGreen, offset: 0, index: 1)
    computeEncoder?.setBuffer(columnDataBlue, offset: 0, index: 2)
    computeEncoder?.setSamplerState(kSamplerState, index: 0)
    computeEncoder?.dispatchThreadgroups(threadgroupCount, threadsPerThreadgroup: threadgroupSize)
    computeEncoder?.endEncoding()

// вернуть true 

}

1 Ответ

0 голосов
/ 11 июня 2018

Так что спасибо Кену Томассу за это;это было просто в конце.

Вы должны установить свойства дескриптора сэмплера перед созданием samplerState. 

let samplerDescriptor = MTLSamplerDescriptor()
  samplerDescriptor.normalizedCoordinates = false
  let  kSamplerState = kDevice.makeSamplerState(descriptor: samplerDescriptor)
...