Cách tăng tốc mã kim loại cho iOS/Mac OS

Question

Tôi đang cố gắng triển khai mã trong Metal thực hiện phép chuyển đổi 1D giữa hai vectơ có độ dài. Tôi đã thực hiện những điều sau đây mà làm việc một cách chính xác

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; 
    } 
} 

Vấn đề của tôi là phải mất lâu hơn khoảng 10 lần để xử lý (truyền dữ liệu tính toán + đến/từ GPU) cùng một dữ liệu sử dụng kim loại hơn trong Swift trên CPU. Câu hỏi của tôi là làm cách nào để thay thế vòng lặp bên trong bằng một thao tác vectơ đơn lẻ hoặc có cách nào khác để tăng tốc mã trên không?

Answer 1

Chìa khóa để lợi dụng tính song song của GPU trong trường hợp này là để nó quản lý vòng lặp ngoài cho bạn. Thay vì gọi hạt nhân một lần cho toàn bộ vectơ dữ liệu, chúng ta sẽ gọi nó cho mỗi phần tử trong vectơ dữ liệu. Chức năng hạt nhân đơn giản hóa điều này:

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; 
} 

Để gửi tác phẩm này, chúng tôi chọn kích thước nhóm dựa trên chiều rộng thực hiện luồng được đề xuất bởi trạng thái đường ống tính toán. Một điều khó khăn ở đây là đảm bảo rằng có đủ đệm trong bộ đệm đầu vào và đầu ra để chúng tôi có thể vượt quá kích thước thực tế của dữ liệu. Điều này làm cho chúng ta lãng phí một lượng nhỏ bộ nhớ và tính toán, nhưng tiết kiệm cho chúng ta sự phức tạp của việc thực hiện một công văn riêng biệt chỉ để tính toán các chập đối với các phần tử ở phần cuối của bộ đệm.

// 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() 

Trong các thí nghiệm của tôi, cách tiếp cận song song này chạy 400-1000x nhanh so với phiên bản nối tiếp trong câu hỏi. Tôi tò mò muốn biết nó so sánh với việc triển khai CPU của bạn như thế nào.