ios - usar - quitar memoria compartida del sistema

Rendimiento de escritura de la memoria: memoria compartida de la CPU GPU

Question

Estoy asignando MTLBuffer entrada y salida usando posix_memalign acuerdo con la documentación compartida de GPU / CPU proporcionada por memkite.

Aparte: es más fácil usar la API más reciente que estiércol con posix_memalign

let metalBuffer = self.metalDevice.newBufferWithLength(byteCount, options: .StorageModeShared)

Mi función kernel opera en aproximadamente 16 millones de estructuras de valores complejos y escribe una cantidad igual de estructuras de valores complejos en la memoria.

Realicé algunos experimentos y mi ''sección matemática compleja'' del kernel de Metal se ejecuta en 0.003 segundos (¡Sí!), Pero escribir el resultado en el buffer lleva> 0.05 (No!) Segundos. En mi experimento, comenté la parte matemática y simplemente asigné el cero a la memoria y demora 0.05 segundos, comentando la tarea y agregando la parte posterior de la matemática, 0.003 segundos.

¿La memoria compartida es lenta en este caso, o hay alguna otra sugerencia o truco que pueda probar?

Detalle adicional

Plataformas de prueba

• iPhone 6S - ~ 0.039 segundos por cuadro
• iPad Air 2 - ~ 0.130 segundos por cuadro

La transmisión de datos

Cada actualización del sombreador recibe aproximadamente 50000 números complejos en forma de un par de tipos de float en una estructura.

struct ComplexNumber {
    float real;
    float imaginary;
};

Firma del kernel

kernel void processChannelData(const device Parameters *parameters [[ buffer(0) ]],
                               const device ComplexNumber *inputSampleData [[ buffer(1) ]],
                               const device ComplexNumber *partAs [[ buffer(2) ]],
                               const device float *partBs [[ buffer(3) ]],
                               const device int *lookups [[ buffer(4) ]],
                               device float *outputImageData [[ buffer(5) ]],
                               uint threadIdentifier [[ thread_position_in_grid ]]);

Todos los búferes contienen, actualmente, datos invariables, excepto inputSampleData que recibe las 50000 muestras en las que inputSampleData . Los otros almacenamientos intermedios contienen aproximadamente 16 millones de valores (128 canales x 130000 píxeles) cada uno. Realizo algunas operaciones en cada ''píxel'' y sumo el resultado complejo a través de los canales y finalmente tomo el valor absoluto del número complejo y asigno el float resultante a outputImageData .

Envío

commandEncoder.setComputePipelineState(pipelineState)

commandEncoder.setBuffer(parametersMetalBuffer, offset: 0, atIndex: 0)
commandEncoder.setBuffer(inputSampleDataMetalBuffer, offset: 0, atIndex: 1)
commandEncoder.setBuffer(partAsMetalBuffer, offset: 0, atIndex: 2)
commandEncoder.setBuffer(partBsMetalBuffer, offset: 0, atIndex: 3)
commandEncoder.setBuffer(lookupsMetalBuffer, offset: 0, atIndex: 4)
commandEncoder.setBuffer(outputImageDataMetalBuffer, offset: 0, atIndex: 5)

let threadExecutionWidth = pipelineState.threadExecutionWidth
let threadsPerThreadgroup = MTLSize(width: threadExecutionWidth, height: 1, depth: 1)
let threadGroups = MTLSize(width: self.numberOfPixels / threadsPerThreadgroup.width, height: 1, depth:1)

commandEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadsPerThreadgroup)
commandEncoder.endEncoding()
metalCommandBuffer.commit()
metalCommandBuffer.waitUntilCompleted()

Ejemplo de GitHub

Escribí un ejemplo llamado Slow y lo puse en GitHub. Parece que el cuello de botella es la escritura de los valores en el Buffer de entrada. Entonces, supongo que la pregunta es ¿cómo evitar el cuello de botella?

Copia de memoria

Escribí una prueba rápida para comparar el rendimiento de varios métodos de copia de bytes.

Estado actual

Reduje el tiempo de ejecución a 0.02 segundos, lo que no parece mucho, pero hace una gran diferencia en el número de fotogramas por segundo. Actualmente, las mejoras más importantes son el resultado de cambiar a cblas_scopy() .

Answer 1

Reducir el tamaño del tipo

Originalmente, estaba pre-convirtiendo los enteros de 16 bits firmados como Floats (32 bits) ya que en última instancia así es como se usarán. Este es un caso donde el rendimiento comienza obligándote a almacenar los valores en 16 bits para reducir el tamaño de tus datos a la mitad.

Objective-C sobre Swift

Para el código que trata sobre el movimiento de datos, puede elegir Objective-C sobre Swift (recomendación de Warren Moore). El rendimiento de Swift en estas situaciones especiales todavía no está a la altura. También puede intentar llamar a memcpy o métodos similares. He visto un par de ejemplos que utilizaron Búfers de bucle for-loop y esto en mis experimentos se realizó lentamente.

Dificultad de prueba

Realmente quería hacer algunos de los experimentos con relación a varios métodos de copia en un patio de recreo en la máquina y desafortunadamente esto era inútil. Las versiones de los dispositivos iOS de los mismos experimentos se realizaron de forma completamente diferente. Uno podría pensar que el rendimiento relativo sería similar, pero descubrí que esto también es una suposición inválida. Sería muy conveniente si pudieras tener un patio de recreo que utilizara el dispositivo iOS como intérprete.

Answer 2

Puede obtener una gran aceleración mediante la codificación de sus datos a los códigos huffman y la decodificación en la GPU, consulte MetalHuffman . Sin embargo, depende de tus datos.