如何加速iOS/Mac OS的Metal代码

Question

如何加速iOS/Mac OS的Metal代码

7

我正在尝试在Metal中实现一段代码，用于执行两个长度向量之间的一维卷积。我已经实现了以下代码，并且它可以正确工作。

kernel void convolve(const device float *dataVector [[ buffer(0) ]],
                     const device int& dataSize [[ buffer(1) ]],
                     const device float *filterVector [[ buffer(2) ]],
                     const device int& filterSize [[ buffer(3) ]],
                     device float *outVector [[ buffer(4) ]],
                     uint id [[ thread_position_in_grid ]]) {
    int outputSize = dataSize - filterSize + 1;
    for (int i=0;i<outputSize;i++) {
        float sum = 0.0;
        for (int j=0;j<filterSize;j++) {
            sum += dataVector[i+j] * filterVector[j];
        }
        outVector[i] = sum;
    }
}

我的问题是通过Metal处理相同数据需要的时间（计算+数据传输到/从GPU）比使用Swift在CPU上处理要慢大约10倍。我的问题是如何用单个向量操作替换内部循环，或者是否有其他方法来加速上述代码？

- Epsilon

1

你的内核函数完全是串行编写的，没有利用GPU的并行性。但在你开始优化之前，你的数据向量有多大，它改变的频率如何？如果数据传输时间占主导地位，使用GPU可能不是正确的方法。 - warrenm

是的，就像@warrenm已经指出的那样，您没有利用GPU的并行性。这不是GPU有效处理事物的方法。您必须将数据发送到GPU，以便每个片段计算单独的乘法范围。 - codetiger

GPU的示例在这里：http://stackoverflow.com/questions/12576976/1d-convolution-without-if-else-statements-non-fft - codetiger

@warrenm 数据向量大约有10,000个（但可能更大），滤波器大小约为64。我的时间测试表明，使用上述代码处理数据比将数据传输到/从GPU花费的时间要长得多。 - Epsilon

@warrenm，你知道如何将上述代码重编以进行并行处理吗？ - Epsilon

2个回答

-1

以下代码展示了如何使用Objective-C Metal API在GPU上并行渲染编码命令（上面的线程代码仅将输出渲染分成网格部分以进行并行处理；计算仍未并行执行）。这就是你在问题中提到的，尽管它不完全符合你的要求。我提供这个答案来帮助任何可能遇到这个问题的人，认为它会提供与并行渲染相关的答案（实际上并不是这样）。

    - (void)drawInMTKView:(MTKView *)view
    {
        dispatch_async(((AppDelegate *)UIApplication.sharedApplication.delegate).cameraViewQueue, ^{
                    id <CAMetalDrawable> drawable = [view currentDrawable]; //[(CAMetalLayer *)view.layer nextDrawable];
                    MTLRenderPassDescriptor *renderPassDesc = [view currentRenderPassDescriptor];
                    renderPassDesc.colorAttachments[0].loadAction = MTLLoadActionClear;
                    renderPassDesc.colorAttachments[0].clearColor = MTLClearColorMake(0.0,0.0,0.0,1.0);
                    renderPassDesc.renderTargetWidth = self.texture.width;
                    renderPassDesc.renderTargetHeight = self.texture.height;
                    renderPassDesc.colorAttachments[0].texture = drawable.texture;
                    if (renderPassDesc != nil)
                    {
                        dispatch_semaphore_wait(self._inflight_semaphore, DISPATCH_TIME_FOREVER);
                        id <MTLCommandBuffer> commandBuffer = [self.metalContext.commandQueue commandBuffer];
                        [commandBuffer enqueue];
            // START PARALLEL RENDERING OPERATIONS HERE
                        id <MTLParallelRenderCommandEncoder> parallelRCE = [commandBuffer parallelRenderCommandEncoderWithDescriptor:renderPassDesc];
// FIRST PARALLEL RENDERING OPERATION
                        id <MTLRenderCommandEncoder> renderEncoder = [parallelRCE renderCommandEncoder];

                        [renderEncoder setRenderPipelineState:self.metalContext.renderPipelineState];

                        [renderEncoder setVertexBuffer:self.metalContext.vertexBuffer offset:0 atIndex:0];
                        [renderEncoder setVertexBuffer:self.metalContext.uniformBuffer offset:0 atIndex:1];

                        [renderEncoder setFragmentBuffer:self.metalContext.uniformBuffer offset:0 atIndex:0];

                        [renderEncoder setFragmentTexture:self.texture
                                                  atIndex:0];

                        [renderEncoder drawPrimitives:MTLPrimitiveTypeTriangleStrip
                                          vertexStart:0
                                          vertexCount:4
                                        instanceCount:1];

                        [renderEncoder endEncoding];
            // ADD SECOND, THIRD, ETC. PARALLEL RENDERING OPERATION HERE
.
.
.
// SUBMIT ALL RENDERING OPERATIONS IN PARALLEL HERE
                        [parallelRCE endEncoding];

                        __block dispatch_semaphore_t block_sema = self._inflight_semaphore;
                        [commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> buffer) {
                            dispatch_semaphore_signal(block_sema);

                        }];

                        if (drawable)
                            [commandBuffer presentDrawable:drawable];
                        [commandBuffer commit];
                        [commandBuffer waitUntilScheduled];
                    }
        });
    }

在上面的例子中，您需要为每个想要并行执行的计算复制与renderEncoder相关的内容。我不认为这对您的代码示例有什么好处，因为一个操作似乎依赖于另一个操作。因此，您最好能得到warrenm提供给您的代码，尽管那并不能真正称为并行渲染。

- James Bush

OP的问题没有提到渲染。并行渲染命令编码是一种专门的技术，只有在命令编码变成CPU瓶颈时才有用。 - warrenm

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- warrenm · Accepted Answer

在这种情况下利用GPU的并行性的关键是让它为您管理外层循环。我们不会一次为整个数据向量调用内核，而是为数据向量中的每个元素调用它。内核函数简化为：

kernel void convolve(const device float *dataVector [[ buffer(0) ]],
                     const constant int &dataSize [[ buffer(1) ]],
                     const constant float *filterVector [[ buffer(2) ]],
                     const constant int &filterSize [[ buffer(3) ]],
                     device float *outVector [[ buffer(4) ]],
                     uint id [[ thread_position_in_grid ]])
{
    float sum = 0.0;
    for (int i = 0; i < filterSize; ++i) {
        sum += dataVector[id + i] * filterVector[i];
    }
    outVector[id] = sum;
}

为了执行这项工作，我们根据计算管线状态建议的线程执行宽度选择线程组大小。这里唯一棘手的问题是确保输入和输出缓冲区中有足够的填充空间，以便我们可以略微超出实际数据大小。这确实会导致我们浪费少量内存和计算资源，但省去了为缓冲区末尾的元素单独计算卷积而进行单独调度的复杂性。

// We should ensure here that the data buffer and output buffer each have a size that is a multiple of
// the compute pipeline's threadExecutionWidth, by padding the amount we allocate for each of them.
// After execution, we ignore the extraneous elements in the output buffer beyond the first (dataCount - filterCount + 1).

let iterationCount = dataCount - filterCount + 1
let threadsPerThreadgroup = MTLSize(width: min(iterationCount, computePipeline.threadExecutionWidth), height: 1, depth: 1)
let threadgroups = (iterationCount + threadsPerThreadgroup.width - 1) / threadsPerThreadgroup.width
let threadgroupsPerGrid = MTLSize(width: threadgroups, height: 1, depth: 1)

let commandEncoder = commandBuffer.computeCommandEncoder()
commandEncoder.setComputePipelineState(computePipeline)
commandEncoder.setBuffer(dataBuffer, offset: 0, at: 0)
commandEncoder.setBytes(&dataCount, length: MemoryLayout<Int>.stride, at: 1)
commandEncoder.setBuffer(filterBuffer, offset: 0, at: 2)
commandEncoder.setBytes(&filterCount, length: MemoryLayout<Int>.stride, at: 3)
commandEncoder.setBuffer(outBuffer, offset: 0, at: 4)
commandEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)
commandEncoder.endEncoding()

在我的实验中，这种并行化方法比问题中的串行版本运行速度快400-1000倍，我很想听听它与你的CPU实现相比如何。